#include "cuda.h"
/*---------------------------------------------------------------------
   ______ ______ _____ ______ _____ 
  |  ____|___  // ____|  ____/ ____|
  | |__     / /| (___ | |__ | |  __ 
  |  __|   / /  \___ \|  __|| | |_ |
  | |____ / /__ ____) | |___| |__| |
  |______/_____|_____/|______\_____|
  GPU-enabled version using CUDA
  Version 1.01

  EZSEG: Routine to segment an image using a two-threshold
  variable-connectivity region growing method utilizing
  GPU acceleration through CUDA.

  void ezseg_cuda(float *IMG, float *SEG, int nt, int np,
                  float thresh1, float thresh2, int nc, int* iters)
 
  INPUT/OUTPUT:
        IMG:    Input image.
        SEG: 
             ON INPUT:
                 Matrix of size (nt,np) which contain
                 1's where there is valid IMG data, and
                 non-zero values for areas with invalid/no IMG data.
             ON OUTPUT:
                 Segmentation map (0:detection ,same as input o.w.).
        nt,np:   Dimensions of image.
        thresh1: Seeding threshold value.
        thresh2: Growing threshold value.
        nc:      # of consecutive pixels needed for connectivity.
        iters:
             ON INPUT: 
                 maximum limit on number of iterations.
             ON OUTPUT: 
                 number of iterations performed.

----------------------------------------------------------------------
 Copyright (c) 2015 Predictive Science Inc.
 
 Permission is hereby granted, free of charge, to any person obtaining
 a copy of this software and associated documentation files 
 (the "Software"), to deal in the Software without restriction, 
 including without limitation the rights to use, copy, modify, merge,
 publish, distribute, sublicense, and/or sell copies of the Software,
 and to permit persons to whom the Software is furnished to do so, 
 subject to the following conditions:
 
 The above copyright notice and this permission notice shall be 
 included in all copies or substantial portions of the Software.
 
 THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, 
 EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF 
 MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND 
 NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS
 BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN
 ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN 
 CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
 SOFTWARE.
 ----------------------------------------------------------------------
*/

/*Define block size.*/
const int BLOCK_SIZE = 16;

/*** Kernel for ezseg iteration ***/

__global__ void ezseg_kernel(float *EUV, float *CHM, float *CHM_TMP, int nt, int np,
                             float thresh1, float thresh2, int nc, int *val_modded)
{
  int fillit,ij,ii,jj;

  int i = blockIdx.x*blockDim.x + threadIdx.x;
  int j = blockIdx.y*blockDim.y + threadIdx.y;
  
  float local_vec[15],tmp_sum;

  if((i>0 && i<nt-1) && (j>0 && j<np-1))
  {
    ij = nt*j+i;  /*Location of i,j in 1D global arrays*/ 

    fillit = 0;   

    if(CHM_TMP[ij] == 1 ){  /*Good data not marked yet*/
      
      if(EUV[ij] <= thresh1){
        fillit = 1;
      }
      else if(EUV[ij] <= thresh2){                             
        local_vec[ 0] = CHM_TMP[nt*(j+1)+(i-1)];
        local_vec[ 1] = CHM_TMP[nt*(j+1)+(i  )];
        local_vec[ 2] = CHM_TMP[nt*(j+1)+(i+1)];
        local_vec[ 3] = CHM_TMP[nt*(j  )+(i+1)];
        local_vec[ 4] = CHM_TMP[nt*(j-1)+(i+1)];
        local_vec[ 5] = CHM_TMP[nt*(j-1)+(i  )];
        local_vec[ 6] = CHM_TMP[nt*(j-1)+(i-1)];
        local_vec[ 7] = CHM_TMP[nt*(j  )+(i-1)];
        local_vec[ 8] = local_vec[0];
        local_vec[ 9] = local_vec[1];
        local_vec[10] = local_vec[2];
        local_vec[11] = local_vec[3];
        local_vec[12] = local_vec[4];
        local_vec[13] = local_vec[5];
        local_vec[14] = local_vec[6];

        for(ii=0;ii<8;ii++){  
          tmp_sum = 0.0f;
          for(jj=0;jj<nc;jj++){
            tmp_sum = tmp_sum + local_vec[ii+jj];
          }
          if(tmp_sum == 0){
            fillit = 1;
            break;
          }
        }
      }  /*euv<thresh2*/

      if (fillit == 1) {
        CHM[ij] = 0.0f;
        if(*val_modded == 0) {
          atomicAdd(val_modded, 1);
        }
      } 
    } /*good data no mark*/
  } /*valid point*/
}

/*********************************************************************/
/*********************************************************************/
/*********************************************************************/

extern "C" void ezseg_cuda(float *EUV, float *CHM, int nt, int np,
                           float thresh1, float thresh2, int nc, int* iters )
{
    int val_modded,max_iters,k;

    /*GPU variables:*/
    float *EUVgpu,*CHMgpu,*CHM_TMPgpu;
    int *val_modded_gpu;

    /*Allocate GPU arrays:*/
    cudaMalloc((void **) &EUVgpu,     sizeof(float)*nt*np);
    cudaMalloc((void **) &CHMgpu,     sizeof(float)*nt*np);
    cudaMalloc((void **) &CHM_TMPgpu, sizeof(float)*nt*np);
    cudaMalloc((void **) &val_modded_gpu, sizeof(int));

    /*Copy euv and chm to GPU*/
    cudaMemcpy(EUVgpu,     EUV, sizeof(float)*nt*np, cudaMemcpyHostToDevice);
    cudaMemcpy(CHMgpu,     CHM, sizeof(float)*nt*np, cudaMemcpyHostToDevice);
    cudaMemcpy(CHM_TMPgpu, CHMgpu, sizeof(float)*nt*np, cudaMemcpyDeviceToDevice);

    /*Set up CUDA grid*/
    dim3 dimBlock(BLOCK_SIZE,BLOCK_SIZE);
    dim3 dimGrid((int)ceil((nt+0.0)/dimBlock.x),
                 (int)ceil((np+0.0)/dimBlock.y));

    /*Start main loop*/ 
    max_iters = *iters;
    *iters = 0;
    for(k=0;k<max_iters;k++){  
      /*Reset val_modded*/
      val_modded = 0;
      cudaMemcpy(val_modded_gpu, &val_modded, sizeof(int), cudaMemcpyHostToDevice);

      /*Perform iteration:*/
      ezseg_kernel<<<dimGrid,dimBlock>>>(EUVgpu,CHMgpu,CHM_TMPgpu,nt,np,thresh1,
                                           thresh2,nc,val_modded_gpu);

      *iters = *iters + 1;

      /*Make sure everything is done*/
      cudaDeviceSynchronize();

      /*Get data mod flag*/        
      cudaMemcpy(&val_modded, val_modded_gpu, sizeof(int), cudaMemcpyDeviceToHost);

      /*If no new CH points, break out of iterations*/
      if(val_modded == 0){
        break;
      }

      /*Reset tmp to be new map iterate for next iteration:*/
      cudaMemcpy(CHM_TMPgpu,CHMgpu,sizeof(float)*nt*np,cudaMemcpyDeviceToDevice);

    }

    /*Copy result from GPU back to CPU*/
    cudaMemcpy(CHM,CHMgpu,sizeof(float)*nt*np,cudaMemcpyDeviceToHost);

    /*Free up GPU memory*/
    cudaFree(EUVgpu);
    cudaFree(CHMgpu);
    cudaFree(CHM_TMPgpu);
    cudaFree(val_modded_gpu);
	
}

	.file	"tmpxft_0020ec8e_00000000-6_ezseg_cuda.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2011:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2011:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z43__device_stub__Z12ezseg_kernelPfS_S_iiffiPiPfS_S_iiffiPi
	.type	_Z43__device_stub__Z12ezseg_kernelPfS_S_iiffiPiPfS_S_iiffiPi, @function
_Z43__device_stub__Z12ezseg_kernelPfS_S_iiffiPiPfS_S_iiffiPi:
.LFB2033:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movq	224(%rsp), %rax
	movq	%rdi, 56(%rsp)
	leaq	80(%rsp), %rdi
	movq	%rsi, 48(%rsp)
	leaq	92(%rsp), %rsi
	movq	%rdx, 40(%rsp)
	leaq	64(%rsp), %rdx
	movl	%ecx, 36(%rsp)
	leaq	72(%rsp), %rcx
	movl	%r8d, 32(%rsp)
	movl	%r9d, 20(%rsp)
	movq	%rax, 8(%rsp)
	movss	%xmm0, 28(%rsp)
	movss	%xmm1, 24(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movl	$1, 88(%rsp)
	movq	%rax, 128(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	36(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 80(%rsp)
	movq	%rax, 92(%rsp)
	movl	$1, 100(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	72(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_Z12ezseg_kernelPfS_S_iiffiPi(%rip), %rdi
	pushq	72(%rsp)
	.cfi_def_cfa_offset 240
	movq	108(%rsp), %rcx
	movl	116(%rsp), %r8d
	movq	96(%rsp), %rsi
	movl	104(%rsp), %edx
	leaq	144(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L2:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2033:
	.size	_Z43__device_stub__Z12ezseg_kernelPfS_S_iiffiPiPfS_S_iiffiPi, .-_Z43__device_stub__Z12ezseg_kernelPfS_S_iiffiPiPfS_S_iiffiPi
	.globl	_Z12ezseg_kernelPfS_S_iiffiPi
	.type	_Z12ezseg_kernelPfS_S_iiffiPi, @function
_Z12ezseg_kernelPfS_S_iiffiPi:
.LFB2034:
	.cfi_startproc
	endbr64
	jmp	_Z43__device_stub__Z12ezseg_kernelPfS_S_iiffiPiPfS_S_iiffiPi
	.cfi_endproc
.LFE2034:
	.size	_Z12ezseg_kernelPfS_S_iiffiPi, .-_Z12ezseg_kernelPfS_S_iiffiPi
	.globl	ezseg_cuda
	.type	ezseg_cuda, @function
ezseg_cuda:
.LFB2008:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movabsq	$68719476752, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movq	%rdi, %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movl	%edx, %ebp
	movslq	%ecx, %rdx
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rdx, %r12
	movq	%r9, %rbx
	subq	$120, %rsp
	.cfi_def_cfa_offset 176
	movl	%r8d, 28(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 8(%rsp)
	movss	%xmm1, 24(%rsp)
	movss	%xmm0, 20(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	movslq	%ebp, %rax
	imulq	%rdx, %rax
	salq	$2, %rax
	movq	%rax, %rsi
	movq	%rax, (%rsp)
	call	cudaMalloc@PLT
	movq	(%rsp), %rsi
	leaq	56(%rsp), %rdi
	call	cudaMalloc@PLT
	movq	(%rsp), %rsi
	leaq	64(%rsp), %rdi
	call	cudaMalloc@PLT
	leaq	72(%rsp), %rdi
	movl	$4, %esi
	call	cudaMalloc@PLT
	movq	(%rsp), %rdx
	movq	48(%rsp), %rdi
	movq	%r13, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movq	(%rsp), %rdx
	movq	8(%rsp), %rsi
	movl	$1, %ecx
	movq	56(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	64(%rsp), %rdi
	movq	(%rsp), %rdx
	movl	$3, %ecx
	movq	56(%rsp), %rsi
	call	cudaMemcpy@PLT
	cvtsi2sdl	%r12d, %xmm0
	mulsd	.LC0(%rip), %xmm0
	movl	$1, 88(%rsp)
	call	ceil@PLT
	movq	%xmm0, %r13
	cvtsi2sdl	%ebp, %xmm0
	mulsd	.LC0(%rip), %xmm0
	call	ceil@PLT
	movq	%r13, %xmm2
	xorl	%edi, %edi
	movl	(%rbx), %r14d
	cvttsd2sil	%xmm0, %eax
	movl	$1, 100(%rsp)
	xorl	%r13d, %r13d
	movl	%edi, (%rbx)
	movl	%eax, 92(%rsp)
	cvttsd2sil	%xmm2, %eax
	movl	%eax, 96(%rsp)
.L9:
	cmpl	%r14d, %r13d
	jge	.L13
	xorl	%esi, %esi
	movq	72(%rsp), %rdi
	movl	$1, %ecx
	movl	$4, %edx
	movl	%esi, 44(%rsp)
	leaq	44(%rsp), %rsi
	call	cudaMemcpy@PLT
	movl	88(%rsp), %ecx
	movl	100(%rsp), %esi
	xorl	%r9d, %r9d
	movq	92(%rsp), %rdi
	xorl	%r8d, %r8d
	movq	%r15, %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L10
	pushq	%rax
	.cfi_def_cfa_offset 184
	movl	%ebp, %ecx
	movl	%r12d, %r8d
	pushq	80(%rsp)
	.cfi_def_cfa_offset 192
	movl	44(%rsp), %r9d
	movss	40(%rsp), %xmm1
	movss	36(%rsp), %xmm0
	movq	80(%rsp), %rdx
	movq	72(%rsp), %rsi
	movq	64(%rsp), %rdi
	call	_Z43__device_stub__Z12ezseg_kernelPfS_S_iiffiPiPfS_S_iiffiPi
	popq	%rdx
	.cfi_def_cfa_offset 184
	popq	%rcx
	.cfi_def_cfa_offset 176
.L10:
	incl	(%rbx)
	call	cudaDeviceSynchronize@PLT
	movq	72(%rsp), %rsi
	leaq	44(%rsp), %rdi
	movl	$2, %ecx
	movl	$4, %edx
	call	cudaMemcpy@PLT
	cmpl	$0, 44(%rsp)
	jne	.L11
.L13:
	movq	(%rsp), %rdx
	movq	56(%rsp), %rsi
	movl	$2, %ecx
	movq	8(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	48(%rsp), %rdi
	call	cudaFree@PLT
	movq	56(%rsp), %rdi
	call	cudaFree@PLT
	movq	64(%rsp), %rdi
	call	cudaFree@PLT
	movq	72(%rsp), %rdi
	call	cudaFree@PLT
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L14
	call	__stack_chk_fail@PLT
.L11:
	movq	(%rsp), %rdx
	movq	56(%rsp), %rsi
	movl	$3, %ecx
	incl	%r13d
	movq	64(%rsp), %rdi
	call	cudaMemcpy@PLT
	jmp	.L9
.L14:
	addq	$120, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2008:
	.size	ezseg_cuda, .-ezseg_cuda
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"_Z12ezseg_kernelPfS_S_iiffiPi"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2036:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC1(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z12ezseg_kernelPfS_S_iiffiPi(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2036:
	.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.cst8,"aM",@progbits,8
	.align 8
.LC0:
	.long	0
	.long	1068498944
	.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:

#include "hip/hip_runtime.h"
/*---------------------------------------------------------------------
   ______ ______ _____ ______ _____ 
  |  ____|___  // ____|  ____/ ____|
  | |__     / /| (___ | |__ | |  __ 
  |  __|   / /  \___ \|  __|| | |_ |
  | |____ / /__ ____) | |___| |__| |
  |______/_____|_____/|______\_____|
  GPU-enabled version using CUDA
  Version 1.01

  EZSEG: Routine to segment an image using a two-threshold
  variable-connectivity region growing method utilizing
  GPU acceleration through CUDA.

  void ezseg_cuda(float *IMG, float *SEG, int nt, int np,
                  float thresh1, float thresh2, int nc, int* iters)
 
  INPUT/OUTPUT:
        IMG:    Input image.
        SEG: 
             ON INPUT:
                 Matrix of size (nt,np) which contain
                 1's where there is valid IMG data, and
                 non-zero values for areas with invalid/no IMG data.
             ON OUTPUT:
                 Segmentation map (0:detection ,same as input o.w.).
        nt,np:   Dimensions of image.
        thresh1: Seeding threshold value.
        thresh2: Growing threshold value.
        nc:      # of consecutive pixels needed for connectivity.
        iters:
             ON INPUT: 
                 maximum limit on number of iterations.
             ON OUTPUT: 
                 number of iterations performed.

----------------------------------------------------------------------
 Copyright (c) 2015 Predictive Science Inc.
 
 Permission is hereby granted, free of charge, to any person obtaining
 a copy of this software and associated documentation files 
 (the "Software"), to deal in the Software without restriction, 
 including without limitation the rights to use, copy, modify, merge,
 publish, distribute, sublicense, and/or sell copies of the Software,
 and to permit persons to whom the Software is furnished to do so, 
 subject to the following conditions:
 
 The above copyright notice and this permission notice shall be 
 included in all copies or substantial portions of the Software.
 
 THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, 
 EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF 
 MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND 
 NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS
 BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN
 ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN 
 CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
 SOFTWARE.
 ----------------------------------------------------------------------
*/

/*Define block size.*/
const int BLOCK_SIZE = 16;

/*** Kernel for ezseg iteration ***/

__global__ void ezseg_kernel(float *EUV, float *CHM, float *CHM_TMP, int nt, int np,
                             float thresh1, float thresh2, int nc, int *val_modded)
{
  int fillit,ij,ii,jj;

  int i = blockIdx.x*blockDim.x + threadIdx.x;
  int j = blockIdx.y*blockDim.y + threadIdx.y;
  
  float local_vec[15],tmp_sum;

  if((i>0 && i<nt-1) && (j>0 && j<np-1))
  {
    ij = nt*j+i;  /*Location of i,j in 1D global arrays*/ 

    fillit = 0;   

    if(CHM_TMP[ij] == 1 ){  /*Good data not marked yet*/
      
      if(EUV[ij] <= thresh1){
        fillit = 1;
      }
      else if(EUV[ij] <= thresh2){                             
        local_vec[ 0] = CHM_TMP[nt*(j+1)+(i-1)];
        local_vec[ 1] = CHM_TMP[nt*(j+1)+(i  )];
        local_vec[ 2] = CHM_TMP[nt*(j+1)+(i+1)];
        local_vec[ 3] = CHM_TMP[nt*(j  )+(i+1)];
        local_vec[ 4] = CHM_TMP[nt*(j-1)+(i+1)];
        local_vec[ 5] = CHM_TMP[nt*(j-1)+(i  )];
        local_vec[ 6] = CHM_TMP[nt*(j-1)+(i-1)];
        local_vec[ 7] = CHM_TMP[nt*(j  )+(i-1)];
        local_vec[ 8] = local_vec[0];
        local_vec[ 9] = local_vec[1];
        local_vec[10] = local_vec[2];
        local_vec[11] = local_vec[3];
        local_vec[12] = local_vec[4];
        local_vec[13] = local_vec[5];
        local_vec[14] = local_vec[6];

        for(ii=0;ii<8;ii++){  
          tmp_sum = 0.0f;
          for(jj=0;jj<nc;jj++){
            tmp_sum = tmp_sum + local_vec[ii+jj];
          }
          if(tmp_sum == 0){
            fillit = 1;
            break;
          }
        }
      }  /*euv<thresh2*/

      if (fillit == 1) {
        CHM[ij] = 0.0f;
        if(*val_modded == 0) {
          atomicAdd(val_modded, 1);
        }
      } 
    } /*good data no mark*/
  } /*valid point*/
}

/*********************************************************************/
/*********************************************************************/
/*********************************************************************/

extern "C" void ezseg_cuda(float *EUV, float *CHM, int nt, int np,
                           float thresh1, float thresh2, int nc, int* iters )
{
    int val_modded,max_iters,k;

    /*GPU variables:*/
    float *EUVgpu,*CHMgpu,*CHM_TMPgpu;
    int *val_modded_gpu;

    /*Allocate GPU arrays:*/
    hipMalloc((void **) &EUVgpu,     sizeof(float)*nt*np);
    hipMalloc((void **) &CHMgpu,     sizeof(float)*nt*np);
    hipMalloc((void **) &CHM_TMPgpu, sizeof(float)*nt*np);
    hipMalloc((void **) &val_modded_gpu, sizeof(int));

    /*Copy euv and chm to GPU*/
    hipMemcpy(EUVgpu,     EUV, sizeof(float)*nt*np, hipMemcpyHostToDevice);
    hipMemcpy(CHMgpu,     CHM, sizeof(float)*nt*np, hipMemcpyHostToDevice);
    hipMemcpy(CHM_TMPgpu, CHMgpu, sizeof(float)*nt*np, hipMemcpyDeviceToDevice);

    /*Set up CUDA grid*/
    dim3 dimBlock(BLOCK_SIZE,BLOCK_SIZE);
    dim3 dimGrid((int)ceil((nt+0.0)/dimBlock.x),
                 (int)ceil((np+0.0)/dimBlock.y));

    /*Start main loop*/ 
    max_iters = *iters;
    *iters = 0;
    for(k=0;k<max_iters;k++){  
      /*Reset val_modded*/
      val_modded = 0;
      hipMemcpy(val_modded_gpu, &val_modded, sizeof(int), hipMemcpyHostToDevice);

      /*Perform iteration:*/
      ezseg_kernel<<<dimGrid,dimBlock>>>(EUVgpu,CHMgpu,CHM_TMPgpu,nt,np,thresh1,
                                           thresh2,nc,val_modded_gpu);

      *iters = *iters + 1;

      /*Make sure everything is done*/
      hipDeviceSynchronize();

      /*Get data mod flag*/        
      hipMemcpy(&val_modded, val_modded_gpu, sizeof(int), hipMemcpyDeviceToHost);

      /*If no new CH points, break out of iterations*/
      if(val_modded == 0){
        break;
      }

      /*Reset tmp to be new map iterate for next iteration:*/
      hipMemcpy(CHM_TMPgpu,CHMgpu,sizeof(float)*nt*np,hipMemcpyDeviceToDevice);

    }

    /*Copy result from GPU back to CPU*/
    hipMemcpy(CHM,CHMgpu,sizeof(float)*nt*np,hipMemcpyDeviceToHost);

    /*Free up GPU memory*/
    hipFree(EUVgpu);
    hipFree(CHMgpu);
    hipFree(CHM_TMPgpu);
    hipFree(val_modded_gpu);
	
}

	.text
	.file	"ezseg_cuda.hip"
	.globl	_Z27__device_stub__ezseg_kernelPfS_S_iiffiPi # -- Begin function _Z27__device_stub__ezseg_kernelPfS_S_iiffiPi
	.type	_Z27__device_stub__ezseg_kernelPfS_S_iiffiPi,@function
_Z27__device_stub__ezseg_kernelPfS_S_iiffiPi: # @_Z27__device_stub__ezseg_kernelPfS_S_iiffiPi
	.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	$176, %rsp
	.cfi_def_cfa_offset 224
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	20(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	16(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	12(%rsp), %r8
	movss	%xmm0, (%r8)
	leaq	8(%rsp), %r10
	movss	%xmm1, (%r10)
	leaq	4(%rsp), %r11
	movl	%r9d, (%r11)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	movq	%r10, 48(%rbx)
	movq	%r11, 56(%rbx)
	leaq	224(%rsp), %rax
	movq	%rax, 64(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	32(%rsp), %r12
	leaq	24(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z12ezseg_kernelPfS_S_iiffiPi, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$192, %rsp
	.cfi_adjust_cfa_offset -192
	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	_Z27__device_stub__ezseg_kernelPfS_S_iiffiPi, .Lfunc_end0-_Z27__device_stub__ezseg_kernelPfS_S_iiffiPi
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function ezseg_cuda
.LCPI1_0:
	.quad	0x3fb0000000000000              # double 0.0625
	.text
	.globl	ezseg_cuda
	.type	ezseg_cuda,@function
ezseg_cuda:                             # @ezseg_cuda
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$104, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%r9, %rbx
	movl	%r8d, 76(%rsp)                  # 4-byte Spill
	movss	%xmm1, 72(%rsp)                 # 4-byte Spill
	movss	%xmm0, 68(%rsp)                 # 4-byte Spill
	movq	%rsi, %r13
	movq	%rdi, %r14
	movl	%edx, 60(%rsp)                  # 4-byte Spill
	movslq	%edx, %rbp
	movl	%ecx, 64(%rsp)                  # 4-byte Spill
	movslq	%ecx, %rax
	movq	%rax, 88(%rsp)                  # 8-byte Spill
	movq	%rbp, %r12
	imulq	%rax, %r12
	shlq	$2, %r12
	leaq	48(%rsp), %rdi
	movq	%r12, %rsi
	callq	hipMalloc
	leaq	40(%rsp), %r15
	movq	%r15, %rdi
	movq	%r12, %rsi
	callq	hipMalloc
	leaq	32(%rsp), %rdi
	movq	%r12, %rsi
	callq	hipMalloc
	leaq	24(%rsp), %rdi
	movl	$4, %esi
	callq	hipMalloc
	leaq	48(%rsp), %rax
	movq	(%rax), %rdi
	movq	%r14, %rsi
	movq	%r12, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r15), %rdi
	movq	%r13, 96(%rsp)                  # 8-byte Spill
	movq	%r13, %rsi
	movq	%r12, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	32(%rsp), %rax
	movq	(%rax), %rdi
	movq	(%r15), %rsi
	movq	%r12, %rdx
	movl	$3, %ecx
	callq	hipMemcpy
	cvtsi2sd	%ebp, %xmm0
	mulsd	.LCPI1_0(%rip), %xmm0
	callq	ceil@PLT
	movsd	%xmm0, 80(%rsp)                 # 8-byte Spill
	xorps	%xmm0, %xmm0
	cvtsi2sdl	88(%rsp), %xmm0         # 4-byte Folded Reload
	mulsd	.LCPI1_0(%rip), %xmm0
	callq	ceil@PLT
	movl	(%rbx), %r13d
	movl	$0, (%rbx)
	testl	%r13d, %r13d
	jle	.LBB1_6
# %bb.1:                                # %.lr.ph.preheader
	cvttsd2si	80(%rsp), %eax          # 8-byte Folded Reload
	cvttsd2si	%xmm0, %r14d
	shlq	$32, %r14
	orq	%rax, %r14
	leaq	20(%rsp), %rbp
	movabsq	$68719476752, %r15              # imm = 0x1000000010
.LBB1_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	$0, 20(%rsp)
	movq	24(%rsp), %rdi
	movl	$4, %edx
	movq	%rbp, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_4
# %bb.3:                                #   in Loop: Header=BB1_2 Depth=1
	movq	48(%rsp), %rdi
	movq	40(%rsp), %rsi
	movq	32(%rsp), %rdx
	movq	24(%rsp), %rax
	movq	%rax, (%rsp)
	movl	60(%rsp), %ecx                  # 4-byte Reload
	movl	64(%rsp), %r8d                  # 4-byte Reload
	movss	68(%rsp), %xmm0                 # 4-byte Reload
                                        # xmm0 = mem[0],zero,zero,zero
	movss	72(%rsp), %xmm1                 # 4-byte Reload
                                        # xmm1 = mem[0],zero,zero,zero
	movl	76(%rsp), %r9d                  # 4-byte Reload
	callq	_Z27__device_stub__ezseg_kernelPfS_S_iiffiPi
.LBB1_4:                                #   in Loop: Header=BB1_2 Depth=1
	incl	(%rbx)
	callq	hipDeviceSynchronize
	movq	24(%rsp), %rsi
	movl	$4, %edx
	movq	%rbp, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	cmpl	$0, 20(%rsp)
	je	.LBB1_6
# %bb.5:                                #   in Loop: Header=BB1_2 Depth=1
	movq	32(%rsp), %rdi
	movq	40(%rsp), %rsi
	movq	%r12, %rdx
	movl	$3, %ecx
	callq	hipMemcpy
	decl	%r13d
	jne	.LBB1_2
.LBB1_6:                                # %._crit_edge
	movq	40(%rsp), %rsi
	movq	96(%rsp), %rdi                  # 8-byte Reload
	movq	%r12, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	48(%rsp), %rdi
	callq	hipFree
	movq	40(%rsp), %rdi
	callq	hipFree
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	addq	$104, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	ezseg_cuda, .Lfunc_end1-ezseg_cuda
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB2_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB2_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z12ezseg_kernelPfS_S_iiffiPi, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end2:
	.size	__hip_module_ctor, .Lfunc_end2-__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	.LBB3_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
.LBB3_2:
	retq
.Lfunc_end3:
	.size	__hip_module_dtor, .Lfunc_end3-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z12ezseg_kernelPfS_S_iiffiPi,@object # @_Z12ezseg_kernelPfS_S_iiffiPi
	.section	.rodata,"a",@progbits
	.globl	_Z12ezseg_kernelPfS_S_iiffiPi
	.p2align	3, 0x0
_Z12ezseg_kernelPfS_S_iiffiPi:
	.quad	_Z27__device_stub__ezseg_kernelPfS_S_iiffiPi
	.size	_Z12ezseg_kernelPfS_S_iiffiPi, 8

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

	.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 _Z27__device_stub__ezseg_kernelPfS_S_iiffiPi
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z12ezseg_kernelPfS_S_iiffiPi
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z12ezseg_kernelPfS_S_iiffiPi ; -- Begin function _Z12ezseg_kernelPfS_S_iiffiPi
	.globl	_Z12ezseg_kernelPfS_S_iiffiPi
	.p2align	8
	.type	_Z12ezseg_kernelPfS_S_iiffiPi,@function
_Z12ezseg_kernelPfS_S_iiffiPi:          ; @_Z12ezseg_kernelPfS_S_iiffiPi
; %bb.0:
	s_load_b32 s2, s[0:1], 0x44
	v_and_b32_e32 v3, 0x3ff, v0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s3, s2, 0xffff
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_mad_u64_u32 v[1:2], null, s14, s3, v[3:4]
	s_mov_b32 s3, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_lt_i32_e32 0, v1
	s_cbranch_execz .LBB0_18
; %bb.1:
	s_load_b128 s[8:11], s[0:1], 0x18
	v_bfe_u32 v0, v0, 10, 10
	s_lshr_b32 s2, s2, 16
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_mad_u64_u32 v[2:3], null, s15, s2, v[0:1]
	s_waitcnt lgkmcnt(0)
	s_add_i32 s2, s8, -1
	s_add_i32 s3, s9, -1
	v_cmp_gt_i32_e32 vcc_lo, s2, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cmp_gt_i32_e64 s2, s3, v2
	v_cmp_lt_i32_e64 s3, 0, v2
	s_and_b32 s2, vcc_lo, s2
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	s_and_b32 s2, s2, s3
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 exec_lo, exec_lo, s2
	s_cbranch_execz .LBB0_18
; %bb.2:
	v_mul_lo_u32 v0, v2, s8
	s_load_b64 s[2:3], s[0:1], 0x10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v2, v0, v1
	v_ashrrev_i32_e32 v3, 31, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[16:17], 2, v[2:3]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v5, vcc_lo, s2, v16
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v6, vcc_lo, s3, v17, vcc_lo
	global_load_b32 v2, v[5:6], off
	s_waitcnt vmcnt(0)
	v_cmp_eq_f32_e32 vcc_lo, 1.0, v2
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_18
; %bb.3:
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b32 s12, s[0:1], 0x28
	s_mov_b32 s9, exec_lo
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s4, v16
	v_add_co_ci_u32_e32 v3, vcc_lo, s5, v17, vcc_lo
	s_mov_b32 s4, -1
	global_load_b32 v2, v[2:3], off
	s_waitcnt vmcnt(0)
	v_cmpx_nge_f32_e32 s10, v2
	s_cbranch_execz .LBB0_14
; %bb.4:
	s_mov_b32 s4, 0
	s_mov_b32 s10, exec_lo
	v_cmpx_ge_f32_e32 s11, v2
	s_cbranch_execz .LBB0_13
; %bb.5:
	v_add_nc_u32_e32 v0, s8, v0
	v_ashrrev_i32_e32 v4, 31, v1
	s_lshl_b32 s4, s8, 1
	s_cmp_gt_i32 s12, 0
	s_mov_b32 s11, 0
	v_ashrrev_i32_e32 v3, 31, v0
	v_add_co_u32 v7, vcc_lo, v0, v1
	v_subrev_nc_u32_e32 v11, s4, v0
	v_add_nc_u32_e32 v2, v0, v1
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v8, vcc_lo, v3, v4, vcc_lo
	s_cselect_b32 s8, -1, 0
	v_add_nc_u32_e32 v9, v11, v1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_ashrrev_i32_e32 v3, 31, v2
	v_lshlrev_b64 v[7:8], 2, v[7:8]
	v_ashrrev_i32_e32 v12, 31, v11
                                        ; implicit-def: $sgpr13
                                        ; implicit-def: $sgpr14
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_ashrrev_i32_e32 v10, 31, v9
	v_lshlrev_b64 v[2:3], 2, v[2:3]
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_u32 v7, vcc_lo, s2, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s3, v8, vcc_lo
	v_add_co_u32 v0, vcc_lo, v11, v1
	v_add_co_ci_u32_e32 v1, vcc_lo, v12, v4, vcc_lo
	v_lshlrev_b64 v[9:10], 2, v[9:10]
	v_add_co_u32 v2, vcc_lo, s2, v2
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_lshlrev_b64 v[0:1], 2, v[0:1]
	v_add_co_ci_u32_e32 v3, vcc_lo, s3, v3, vcc_lo
	v_add_co_u32 v9, vcc_lo, s2, v9
	v_add_co_ci_u32_e32 v10, vcc_lo, s3, v10, vcc_lo
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_u32 v11, vcc_lo, s2, v0
	v_add_co_ci_u32_e32 v12, vcc_lo, s3, v1, vcc_lo
	s_clause 0x5
	global_load_b32 v0, v[7:8], off offset:-4
	global_load_b64 v[1:2], v[2:3], off
	global_load_b64 v[14:15], v[9:10], off offset:-4
	global_load_b32 v4, v[11:12], off offset:4
	global_load_b32 v3, v[5:6], off offset:4
	global_load_b32 v7, v[5:6], off offset:-4
	s_mov_b64 s[2:3], 0
	s_waitcnt vmcnt(4)
	v_dual_mov_b32 v8, v0 :: v_dual_mov_b32 v9, v1
	s_waitcnt vmcnt(3)
	v_dual_mov_b32 v10, v2 :: v_dual_mov_b32 v5, v15
	s_waitcnt vmcnt(1)
	v_dual_mov_b32 v11, v3 :: v_dual_mov_b32 v6, v14
	v_dual_mov_b32 v13, v15 :: v_dual_mov_b32 v12, v4
.LBB0_6:                                ; %.preheader
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_8 Depth 2
	v_mov_b32_e32 v15, 0
	s_and_not1_b32 vcc_lo, exec_lo, s8
	s_cbranch_vccnz .LBB0_9
; %bb.7:                                ; %.lr.ph.preheader
                                        ;   in Loop: Header=BB0_6 Depth=1
	s_mov_b64 s[4:5], s[2:3]
	s_mov_b32 s15, s12
.LBB0_8:                                ; %.lr.ph
                                        ;   Parent Loop BB0_6 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_mov_b32 m0, s4
	s_add_i32 s15, s15, -1
	s_waitcnt vmcnt(0)
	v_movrels_b32_e32 v18, v0
	s_add_u32 s4, s4, 1
	s_addc_u32 s5, s5, 0
	s_cmp_lg_u32 s15, 0
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f32_e32 v15, v15, v18
	s_cbranch_scc1 .LBB0_8
.LBB0_9:                                ; %._crit_edge
                                        ;   in Loop: Header=BB0_6 Depth=1
	s_mov_b32 s5, -1
	s_or_b32 s14, s14, exec_lo
	s_mov_b32 s4, exec_lo
	v_cmpx_neq_f32_e32 0, v15
; %bb.10:                               ;   in Loop: Header=BB0_6 Depth=1
	s_add_u32 s2, s2, 1
	s_addc_u32 s3, s3, 0
	s_cmp_eq_u32 s2, 8
	s_cselect_b32 s5, -1, 0
	s_and_not1_b32 s14, s14, exec_lo
	s_or_not1_b32 s5, s5, exec_lo
; %bb.11:                               ; %Flow115
                                        ;   in Loop: Header=BB0_6 Depth=1
	s_or_b32 exec_lo, exec_lo, s4
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s4, exec_lo, s5
	s_or_b32 s11, s4, s11
	s_and_not1_b32 s4, s13, exec_lo
	s_and_b32 s5, s14, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s13, s4, s5
	s_and_not1_b32 exec_lo, exec_lo, s11
	s_cbranch_execnz .LBB0_6
; %bb.12:                               ; %loop.exit.guard
	s_or_b32 exec_lo, exec_lo, s11
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 s4, s13, exec_lo
.LBB0_13:                               ; %Flow117
	s_or_b32 exec_lo, exec_lo, s10
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_not1_b32 s4, s4, exec_lo
.LBB0_14:                               ; %Flow116
	s_or_b32 exec_lo, exec_lo, s9
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 exec_lo, exec_lo, s4
	s_cbranch_execz .LBB0_18
; %bb.15:                               ; %.loopexit
	s_load_b64 s[0:1], s[0:1], 0x30
	v_add_co_u32 v0, vcc_lo, s6, v16
	v_add_co_ci_u32_e32 v1, vcc_lo, s7, v17, vcc_lo
	v_mov_b32_e32 v2, 0
	global_store_b32 v[0:1], v2, off
	s_waitcnt lgkmcnt(0)
	s_load_b32 s2, s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_cmp_lg_u32 s2, 0
	s_cbranch_scc1 .LBB0_18
; %bb.16:
	s_mov_b32 s2, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mbcnt_lo_u32_b32 v0, s2, 0
	v_cmp_eq_u32_e32 vcc_lo, 0, v0
	s_and_b32 s3, exec_lo, vcc_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 exec_lo, s3
	s_cbranch_execz .LBB0_18
; %bb.17:
	s_bcnt1_i32_b32 s2, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s2
	global_atomic_add_u32 v0, v1, s[0:1]
.LBB0_18:                               ; %.critedge
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z12ezseg_kernelPfS_S_iiffiPi
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 312
		.amdhsa_user_sgpr_count 14
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 19
		.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	_Z12ezseg_kernelPfS_S_iiffiPi, .Lfunc_end0-_Z12ezseg_kernelPfS_S_iiffiPi
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 900
; NumSgprs: 18
; NumVgprs: 19
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 19
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         28
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     by_value
      - .offset:         36
        .size:           4
        .value_kind:     by_value
      - .offset:         40
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         48
        .size:           8
        .value_kind:     global_buffer
      - .offset:         56
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         60
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         64
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         68
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         70
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         72
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         74
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         76
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         78
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         104
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         112
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         120
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 312
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z12ezseg_kernelPfS_S_iiffiPi
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z12ezseg_kernelPfS_S_iiffiPi.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     19
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
