/*

Cuda implementation of the VAI forward solver 

Optimization 2-b: Here we switched to float4 and int4 data types 
                  for more effecient memory access and computing indices.
		  Also, loop unrolling in computing the neighboring elements 
		  contribution is dones is done

Optimization 3  : Since we only use two or few current sources, here I 
                  replaced the sources array by passing the sources as parameters 
		  or copy them to constant memory. Whether there is a current source at
		  an element or not is encoded as the first digit in the first neigbing 
		  idices list
 
		  The redundant elements array is removed 
		  GPU memory usage is about 700MB

Adnan Salman: 11/1/2011

*/

#include <stdio.h>
#include <assert.h>
#include <float.h>
#include <unistd.h>


#define THREADS_PER_BLOCK 32
#define HOST_NAME_SIZE    200

// reduction function declaration (implemented in reduction.cu)
float full_reduce(float *d_input, float *d_output, unsigned int n, 
		  unsigned int maxThreads, unsigned int maxBlocks);


__device__ float DotProdVec4(float4 A, float4 B){
  return (A.x*B.x + A.y*B.y + A.z*B.z + A.w*B.w);
}

__device__ float ComputeFF(int idx, int neighborElemIdx, float4 *d4_v1, int inElementsSize, 
			   float dvv1, float4 *d4_AFF, float Tss, float du3, int j){

    float ss = 0,  uvv;

    if (neighborElemIdx != inElementsSize) {

      int    idx0  = idx*16+2*j;
      float4 ndv11 = d4_v1[neighborElemIdx*2];
      float4 ndv12 = d4_v1[neighborElemIdx*2 + 1];
      float ndv[8] = {ndv11.x, ndv11.y, ndv11.z, ndv11.w, 
		      ndv12.x, ndv12.y, ndv12.z, ndv12.w};
    
      uvv = (dvv1 + ndv[7-j])/2.0;
      ss = DotProdVec4(d4_AFF[idx0], ndv11) + DotProdVec4(d4_AFF[idx0+1], ndv12);

    }
    else uvv = dvv1;
    
    return (Tss * uvv +  du3 - ss);
}

__global__ void  DevComputeff4(int inElementsSize, int4 *d4_IJNZ, float4 *d4_srcs,
			       float4 *d4_AFF, float4 *d4_v1, float4 *ff4, float Tss){

  int idx  = blockIdx.x * blockDim.x + threadIdx.x;  // element index
  int idx0 = idx*2;                                  // type4 data index 

  if (idx <inElementsSize){
    int4 dijnz1 = d4_IJNZ[idx0];                    
    float4 dv11 = d4_v1[idx0];
    float4 du31 = make_float4(0,0,0,0);
    float4 du32 = make_float4(0,0,0,0);

    int srcid = dijnz1.x % 10;
    if (srcid){
      du31 = d4_srcs[(srcid-1)*2];
      du32 = d4_srcs[(srcid-1)*2+1];
    }

    float4 ff1;

    //make sure all argument are the same
    ff1.x = ComputeFF(idx, dijnz1.x/10, d4_v1, inElementsSize, dv11.x, d4_AFF, Tss, du31.x, 0);
    ff1.y = ComputeFF(idx, dijnz1.y, d4_v1, inElementsSize, dv11.y, d4_AFF, Tss, du31.y, 1);
    ff1.z = ComputeFF(idx, dijnz1.z, d4_v1, inElementsSize, dv11.z, d4_AFF, Tss, du31.z, 2);
    ff1.w = ComputeFF(idx, dijnz1.w, d4_v1, inElementsSize, dv11.w, d4_AFF, Tss, du31.w, 3);
    ff4[idx0] = ff1;

    dijnz1 = d4_IJNZ[idx0+1];
    dv11   = d4_v1[idx0+1];
    ff1.x = ComputeFF(idx, dijnz1.x, d4_v1, inElementsSize, dv11.x, d4_AFF, Tss, du32.x, 4);
    ff1.y = ComputeFF(idx, dijnz1.y, d4_v1, inElementsSize, dv11.y, d4_AFF, Tss, du32.y, 5);
    ff1.z = ComputeFF(idx, dijnz1.z, d4_v1, inElementsSize, dv11.z, d4_AFF, Tss, du32.z, 6);
    ff1.w = ComputeFF(idx, dijnz1.w, d4_v1, inElementsSize, dv11.w, d4_AFF, Tss, du32.w, 7);
    ff4[idx0+1] = ff1;

  }
}

__global__ void  UpdateSolution(int inElementsSize, float4 *d4_AL, float4 *d4_v1, float4 *ff4, float4 *diff4_v1){

  //TODO: use the temproray ff4 array to hold the temprary diff4_v1 data 
  int idx = blockIdx.x * blockDim.x + threadIdx.x;

  float4 dff1 = ff4[idx*2];
  float4 dff2 = ff4[idx*2+1];
  float4 dvvv;
  int idx0 = idx*16;

  float4 diffv1 = d4_v1[idx*2];

  dvvv.x = DotProdVec4(d4_AL[idx0],   dff1) + DotProdVec4(d4_AL[idx0+1], dff2);
  dvvv.y = DotProdVec4(d4_AL[idx0+2], dff1) + DotProdVec4(d4_AL[idx0+2+1], dff2);
  dvvv.z = DotProdVec4(d4_AL[idx0+4], dff1) + DotProdVec4(d4_AL[idx0+4+1], dff2);
  dvvv.w = DotProdVec4(d4_AL[idx0+6], dff1) + DotProdVec4(d4_AL[idx0+6+1], dff2);
  d4_v1[idx*2] = dvvv;

  diffv1.x = fabs(diffv1.x - dvvv.x);
  diffv1.y = fabs(diffv1.y - dvvv.y);
  diffv1.z = fabs(diffv1.z - dvvv.z);
  diffv1.w = fabs(diffv1.w - dvvv.w);
  diff4_v1[idx*2] = diffv1;

  dvvv.x = DotProdVec4(d4_AL[idx0+8],  dff1) + DotProdVec4(d4_AL[idx0+8+1], dff2);
  dvvv.y = DotProdVec4(d4_AL[idx0+10], dff1) + DotProdVec4(d4_AL[idx0+10+1], dff2);
  dvvv.z = DotProdVec4(d4_AL[idx0+12], dff1) + DotProdVec4(d4_AL[idx0+12+1], dff2);
  dvvv.w = DotProdVec4(d4_AL[idx0+14], dff1) + DotProdVec4(d4_AL[idx0+14+1], dff2);

  diffv1 = d4_v1[idx*2+1];
  d4_v1[idx*2+1] = dvvv;

  diffv1.x = fabs(diffv1.x - dvvv.x);
  diffv1.y = fabs(diffv1.y - dvvv.y);
  diffv1.z = fabs(diffv1.z - dvvv.z);
  diffv1.w = fabs(diffv1.w - dvvv.w);
  diff4_v1[idx*2+1] = diffv1;

}

int InitializeSolver( int device){

  //TODO: move all cuda initialization here 
  
  static int setdev = 0;

  int deviceCount;
  cudaGetDeviceCount(&deviceCount);
  if (deviceCount == 0){
    fprintf(stderr, "There is no device supporting CUDA ");
    return 1;
  }
  
  if (device < 0 || device >= deviceCount){
    fprintf(stderr, "No such cuda device = %d \n", device);
    return 1;
  }

  else {
    if (!setdev){
      if( cudaSetDevice( device ) != cudaSuccess) {
	fprintf(stderr, "Failed setting cuda device \n ");
	return 1;
      }
      setdev = 1;
    }
  }

  char host[HOST_NAME_SIZE];
  gethostname(host, HOST_NAME_SIZE);

  printf("[%s] using id=%d cuda device.\n", host, device);

  return 0; 

}

extern "C" int SolveVaiCuda(int inElementsSize, int* IJNZ, float *v1, float *AFF, float *AL,
			    int device, float Tss, float *srcs, int num_srcs, float scaledTol, 
			    int maxNumIterations, int checkMax, float eps, int printFlag){

  static int setdev = 0;
  if (setdev == 0){
    InitializeSolver(device);
    setdev = 1;
  }

  int   iterations, check = 0; 
  float loop_step_diff = FLT_MAX, diff = FLT_MAX;

  dim3 blocks(THREADS_PER_BLOCK);
  dim3 grid(inElementsSize/blocks.x + 1);

  float4 *d4_v1, *d4_AL, *d4_AFF, *d4_ff, *d4_srcs, *diff4_v1;
  int4   *d4_IJNZ;

  float * ff = (float *) malloc(inElementsSize*8*sizeof(float));

  assert( cudaMalloc( (void**) &d4_IJNZ, inElementsSize*2*sizeof(int4))   == cudaSuccess);
  assert( cudaMemcpy( d4_IJNZ, IJNZ, inElementsSize*8*sizeof(int), cudaMemcpyHostToDevice) == cudaSuccess);

  assert( cudaMalloc( (void**) &d4_AFF, inElementsSize*16*sizeof(float4))   == cudaSuccess);
  assert( cudaMemcpy( d4_AFF, AFF, inElementsSize*64*sizeof(float), cudaMemcpyHostToDevice) == cudaSuccess);

  assert( cudaMalloc( (void**) &d4_AL, inElementsSize*16*sizeof(float4))   == cudaSuccess);
  assert( cudaMemcpy( d4_AL, AL, inElementsSize*64*sizeof(float), cudaMemcpyHostToDevice) == cudaSuccess);

  assert( cudaMalloc( (void**) &d4_ff, inElementsSize*2*sizeof(float4))   == cudaSuccess);

  assert( cudaMalloc( (void**) &d4_srcs, num_srcs*2*sizeof(float4))   == cudaSuccess);
  assert( cudaMemcpy( d4_srcs, srcs, num_srcs*8*sizeof(float), cudaMemcpyHostToDevice) == cudaSuccess);

  assert( cudaMalloc( (void**) &d4_v1, inElementsSize*2*sizeof(float4))   == cudaSuccess);
  assert( cudaMemset( d4_v1, 0, inElementsSize*2*sizeof(float4)) == cudaSuccess);

  assert( cudaMalloc( (void**) &diff4_v1, inElementsSize*2*sizeof(float4))   == cudaSuccess);
  assert( cudaMemset( diff4_v1, -1, inElementsSize*8*sizeof(float)) == cudaSuccess);

  ///temp stuff
  assert( cudaMemset( d4_ff, 0,  inElementsSize*8*sizeof(float)) == cudaSuccess);
  
  for(iterations = 0; iterations < maxNumIterations && loop_step_diff > scaledTol && 
	check < checkMax; iterations++) {

    DevComputeff4<<<grid, blocks>>>(inElementsSize, d4_IJNZ, d4_srcs, d4_AFF, d4_v1, d4_ff, Tss);
    UpdateSolution<<<grid, blocks>>> (inElementsSize, d4_AL, d4_v1, d4_ff, diff4_v1);

    loop_step_diff = full_reduce((float*)diff4_v1, (float*) diff4_v1, inElementsSize*8, 512, 256);
    if (fabs(loop_step_diff - diff) < eps ) check++;
    else {
      check = 0;
      diff = loop_step_diff;
    }
  }

  assert(cudaMemcpy(v1, d4_v1, (inElementsSize)*2*sizeof(float4), cudaMemcpyDeviceToHost) == cudaSuccess);

  /*
  cudaError_t err = cudaMemcpy(v1, d4_v1, (inElementsSize)*2*sizeof(float4), cudaMemcpyDeviceToHost);

  if (err ==     cudaSuccess)
    printf("Success ");
  else if (err == cudaErrorInvalidValue)
    printf("cudaErrorInvalidValue ");
  else if (err == cudaErrorInvalidDevicePointer)
    printf("cudaErrorInvalidDevicePointer ");
  else if (err == cudaErrorInvalidMemcpyDirection)
    printf("cudaErrorInvalidMemcpyDirection ");
  else 
    printf("Unknown error ");
  */


  cudaThreadSynchronize();

  cudaFree(d4_IJNZ);
  cudaFree(d4_AFF);
  cudaFree(d4_AL);
  cudaFree(d4_ff);
  cudaFree(d4_srcs);
  cudaFree(diff4_v1);
  cudaFree(d4_v1);

  return iterations;
  
}

	.file	"tmpxft_00372044_00000000-6_poisson_vai_cuda.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2046:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2046:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z11DotProdVec46float4S_
	.type	_Z11DotProdVec46float4S_, @function
_Z11DotProdVec46float4S_:
.LFB2040:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2040:
	.size	_Z11DotProdVec46float4S_, .-_Z11DotProdVec46float4S_
	.globl	_Z9ComputeFFiiP6float4ifS0_ffi
	.type	_Z9ComputeFFiiP6float4ifS0_ffi, @function
_Z9ComputeFFiiP6float4ifS0_ffi:
.LFB2041:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2041:
	.size	_Z9ComputeFFiiP6float4ifS0_ffi, .-_Z9ComputeFFiiP6float4ifS0_ffi
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"There is no device supporting CUDA "
.LC1:
	.string	"No such cuda device = %d \n"
.LC2:
	.string	"Failed setting cuda device \n "
.LC3:
	.string	"[%s] using id=%d cuda device.\n"
	.text
	.globl	_Z16InitializeSolveri
	.type	_Z16InitializeSolveri, @function
_Z16InitializeSolveri:
.LFB2042:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movl	%edi, %ebx
	subq	$232, %rsp
	.cfi_def_cfa_offset 256
	movq	%fs:40, %rax
	movq	%rax, 216(%rsp)
	xorl	%eax, %eax
	leaq	12(%rsp), %rdi
	call	cudaGetDeviceCount@PLT
	movl	12(%rsp), %eax
	testl	%eax, %eax
	jne	.L7
	movq	stderr(%rip), %rdi
	leaq	.LC0(%rip), %rdx
	movl	$2, %esi
	jmp	.L17
.L7:
	cmpl	%ebx, %eax
	jle	.L15
	testl	%ebx, %ebx
	jns	.L9
.L15:
	movq	stderr(%rip), %rdi
	movl	%ebx, %ecx
	movl	$2, %esi
	xorl	%eax, %eax
	leaq	.LC1(%rip), %rdx
	call	__fprintf_chk@PLT
.L8:
	movl	$1, %eax
	jmp	.L6
.L9:
	cmpl	$0, _ZZ16InitializeSolveriE6setdev(%rip)
	jne	.L12
	movl	%ebx, %edi
	call	cudaSetDevice@PLT
	testl	%eax, %eax
	je	.L13
	movq	stderr(%rip), %rdi
	leaq	.LC2(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
.L17:
	call	__fprintf_chk@PLT
	jmp	.L8
.L13:
	movl	$1, _ZZ16InitializeSolveriE6setdev(%rip)
.L12:
	leaq	16(%rsp), %rbp
	movl	$200, %edx
	movl	$200, %esi
	movq	%rbp, %rdi
	call	__gethostname_chk@PLT
	movl	%ebx, %ecx
	movq	%rbp, %rdx
	movl	$2, %edi
	leaq	.LC3(%rip), %rsi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	xorl	%eax, %eax
.L6:
	movq	216(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L14
	call	__stack_chk_fail@PLT
.L14:
	addq	$232, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2042:
	.size	_Z16InitializeSolveri, .-_Z16InitializeSolveri
	.globl	_Z56__device_stub__Z13DevComputeff4iP4int4P6float4S2_S2_S2_fiP4int4P6float4S2_S2_S2_f
	.type	_Z56__device_stub__Z13DevComputeff4iP4int4P6float4S2_S2_S2_fiP4int4P6float4S2_S2_S2_f, @function
_Z56__device_stub__Z13DevComputeff4iP4int4P6float4S2_S2_S2_fiP4int4P6float4S2_S2_S2_f:
.LFB2068:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movl	%edi, 44(%rsp)
	leaq	64(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	48(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	56(%rsp), %rcx
	movq	%r8, 8(%rsp)
	movq	%r9, (%rsp)
	movss	%xmm0, 40(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	44(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 112(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 144(%rsp)
	movq	%rsp, %rax
	movq	%rax, 152(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L19
	pushq	56(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z13DevComputeff4iP4int4P6float4S2_S2_S2_f(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 208
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	128(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L19:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L21
	call	__stack_chk_fail@PLT
.L21:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2068:
	.size	_Z56__device_stub__Z13DevComputeff4iP4int4P6float4S2_S2_S2_fiP4int4P6float4S2_S2_S2_f, .-_Z56__device_stub__Z13DevComputeff4iP4int4P6float4S2_S2_S2_fiP4int4P6float4S2_S2_S2_f
	.globl	_Z13DevComputeff4iP4int4P6float4S2_S2_S2_f
	.type	_Z13DevComputeff4iP4int4P6float4S2_S2_S2_f, @function
_Z13DevComputeff4iP4int4P6float4S2_S2_S2_f:
.LFB2069:
	.cfi_startproc
	endbr64
	jmp	_Z56__device_stub__Z13DevComputeff4iP4int4P6float4S2_S2_S2_fiP4int4P6float4S2_S2_S2_f
	.cfi_endproc
.LFE2069:
	.size	_Z13DevComputeff4iP4int4P6float4S2_S2_S2_f, .-_Z13DevComputeff4iP4int4P6float4S2_S2_S2_f
	.globl	_Z50__device_stub__Z14UpdateSolutioniP6float4S0_S0_S0_iP6float4S0_S0_S0_
	.type	_Z50__device_stub__Z14UpdateSolutioniP6float4S0_S0_S0_iP6float4S0_S0_S0_, @function
_Z50__device_stub__Z14UpdateSolutioniP6float4S0_S0_S0_iP6float4S0_S0_S0_:
.LFB2070:
	.cfi_startproc
	endbr64
	subq	$168, %rsp
	.cfi_def_cfa_offset 176
	movl	%edi, 44(%rsp)
	leaq	64(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	48(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	56(%rsp), %rcx
	movq	%r8, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leaq	44(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 112(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 144(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L24
	pushq	56(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z14UpdateSolutioniP6float4S0_S0_S0_(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 192
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	128(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 184
	popq	%rdx
	.cfi_def_cfa_offset 176
.L24:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L26
	call	__stack_chk_fail@PLT
.L26:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2070:
	.size	_Z50__device_stub__Z14UpdateSolutioniP6float4S0_S0_S0_iP6float4S0_S0_S0_, .-_Z50__device_stub__Z14UpdateSolutioniP6float4S0_S0_S0_iP6float4S0_S0_S0_
	.globl	_Z14UpdateSolutioniP6float4S0_S0_S0_
	.type	_Z14UpdateSolutioniP6float4S0_S0_S0_, @function
_Z14UpdateSolutioniP6float4S0_S0_S0_:
.LFB2071:
	.cfi_startproc
	endbr64
	jmp	_Z50__device_stub__Z14UpdateSolutioniP6float4S0_S0_S0_iP6float4S0_S0_S0_
	.cfi_endproc
.LFE2071:
	.size	_Z14UpdateSolutioniP6float4S0_S0_S0_, .-_Z14UpdateSolutioniP6float4S0_S0_S0_
	.globl	SolveVaiCuda
	.type	SolveVaiCuda, @function
SolveVaiCuda:
.LFB2043:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movd	%xmm0, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movl	%edi, %ebx
	subq	$48, %rsp
	.cfi_def_cfa_offset 96
	cmpl	$0, _ZZ12SolveVaiCudaE6setdev(%rip)
	movss	%xmm1, 12(%rsp)
	movss	%xmm2, 8(%rsp)
	jne	.L30
	movl	%r9d, %edi
	call	_Z16InitializeSolveri
	movl	$1, _ZZ12SolveVaiCudaE6setdev(%rip)
.L30:
	movl	%ebx, %eax
	movss	.LC4(%rip), %xmm4
	xorl	%ebp, %ebp
	xorl	%r12d, %r12d
	shrl	$5, %eax
	leal	0(,%rbx,8), %r13d
	movabsq	$4294967297, %rdx
	incl	%eax
	movq	%rdx, 28(%rsp)
	movaps	%xmm4, %xmm0
	movl	%eax, 36(%rsp)
	movq	%rdx, 40(%rsp)
	movss	%xmm4, 4(%rsp)
.L31:
	cmpl	112(%rsp), %r12d
	jge	.L39
	cmpl	120(%rsp), %ebp
	jge	.L39
	comiss	12(%rsp), %xmm0
	jbe	.L39
	movl	32(%rsp), %ecx
	movq	36(%rsp), %rdi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$32, 24(%rsp)
	movl	44(%rsp), %esi
	movq	24(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L32
	movd	%r14d, %xmm0
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	xorl	%ecx, %ecx
	xorl	%edx, %edx
	xorl	%esi, %esi
	movl	%ebx, %edi
	call	_Z56__device_stub__Z13DevComputeff4iP4int4P6float4S2_S2_S2_fiP4int4P6float4S2_S2_S2_f
.L32:
	movl	32(%rsp), %ecx
	movq	24(%rsp), %rdx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	36(%rsp), %rdi
	movl	44(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L33
	xorl	%r8d, %r8d
	xorl	%ecx, %ecx
	xorl	%edx, %edx
	xorl	%esi, %esi
	movl	%ebx, %edi
	call	_Z50__device_stub__Z14UpdateSolutioniP6float4S0_S0_S0_iP6float4S0_S0_S0_
.L33:
	xorl	%edi, %edi
	movl	$256, %r8d
	movl	$512, %ecx
	movl	%r13d, %edx
	movq	%rdi, %rsi
	incl	%ebp
	call	_Z11full_reducePfS_jjj@PLT
	movss	8(%rsp), %xmm7
	movss	4(%rsp), %xmm3
	xorl	%eax, %eax
	movaps	%xmm0, %xmm1
	subss	4(%rsp), %xmm1
	andps	.LC5(%rip), %xmm1
	movaps	%xmm1, %xmm2
	comiss	%xmm1, %xmm7
	cmpltss	8(%rsp), %xmm2
	cmovbe	%eax, %ebp
	incl	%r12d
	andps	%xmm2, %xmm3
	andnps	%xmm0, %xmm2
	orps	%xmm3, %xmm2
	movss	%xmm2, 4(%rsp)
	jmp	.L31
.L39:
	call	cudaThreadSynchronize@PLT
	xorl	%edi, %edi
	call	cudaFree@PLT
	xorl	%edi, %edi
	call	cudaFree@PLT
	xorl	%edi, %edi
	call	cudaFree@PLT
	xorl	%edi, %edi
	call	cudaFree@PLT
	xorl	%edi, %edi
	call	cudaFree@PLT
	xorl	%edi, %edi
	call	cudaFree@PLT
	xorl	%edi, %edi
	call	cudaFree@PLT
	addq	$48, %rsp
	.cfi_def_cfa_offset 48
	movl	%r12d, %eax
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2043:
	.size	SolveVaiCuda, .-SolveVaiCuda
	.section	.rodata.str1.1
.LC6:
	.string	"_Z14UpdateSolutioniP6float4S0_S0_S0_"
.LC7:
	.string	"_Z13DevComputeff4iP4int4P6float4S2_S2_S2_f"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2073:
	.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	.LC6(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z14UpdateSolutioniP6float4S0_S0_S0_(%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	.LC7(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z13DevComputeff4iP4int4P6float4S2_S2_S2_f(%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
.LFE2073:
	.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
	.local	_ZZ12SolveVaiCudaE6setdev
	.comm	_ZZ12SolveVaiCudaE6setdev,4,4
	.local	_ZZ16InitializeSolveriE6setdev
	.comm	_ZZ16InitializeSolveriE6setdev,4,4
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC4:
	.long	2139095039
	.section	.rodata.cst16,"aM",@progbits,16
	.align 16
.LC5:
	.long	2147483647
	.long	0
	.long	0
	.long	0
	.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:

/*

Cuda implementation of the VAI forward solver 

Optimization 2-b: Here we switched to float4 and int4 data types 
                  for more effecient memory access and computing indices.
		  Also, loop unrolling in computing the neighboring elements 
		  contribution is dones is done

Optimization 3  : Since we only use two or few current sources, here I 
                  replaced the sources array by passing the sources as parameters 
		  or copy them to constant memory. Whether there is a current source at
		  an element or not is encoded as the first digit in the first neigbing 
		  idices list
 
		  The redundant elements array is removed 
		  GPU memory usage is about 700MB

Adnan Salman: 11/1/2011

*/


#include <hip/hip_runtime.h>
#include <stdio.h>
#include <assert.h>
#include <float.h>
#include <unistd.h>


#define THREADS_PER_BLOCK 32
#define HOST_NAME_SIZE    200

// reduction function declaration (implemented in reduction.cu)
float full_reduce(float *d_input, float *d_output, unsigned int n, 
		  unsigned int maxThreads, unsigned int maxBlocks);


__device__ float DotProdVec4(float4 A, float4 B){
  return (A.x*B.x + A.y*B.y + A.z*B.z + A.w*B.w);
}

__device__ float ComputeFF(int idx, int neighborElemIdx, float4 *d4_v1, int inElementsSize, 
			   float dvv1, float4 *d4_AFF, float Tss, float du3, int j){

    float ss = 0,  uvv;

    if (neighborElemIdx != inElementsSize) {

      int    idx0  = idx*16+2*j;
      float4 ndv11 = d4_v1[neighborElemIdx*2];
      float4 ndv12 = d4_v1[neighborElemIdx*2 + 1];
      float ndv[8] = {ndv11.x, ndv11.y, ndv11.z, ndv11.w, 
		      ndv12.x, ndv12.y, ndv12.z, ndv12.w};
    
      uvv = (dvv1 + ndv[7-j])/2.0;
      ss = DotProdVec4(d4_AFF[idx0], ndv11) + DotProdVec4(d4_AFF[idx0+1], ndv12);

    }
    else uvv = dvv1;
    
    return (Tss * uvv +  du3 - ss);
}

__global__ void  DevComputeff4(int inElementsSize, int4 *d4_IJNZ, float4 *d4_srcs,
			       float4 *d4_AFF, float4 *d4_v1, float4 *ff4, float Tss){

  int idx  = blockIdx.x * blockDim.x + threadIdx.x;  // element index
  int idx0 = idx*2;                                  // type4 data index 

  if (idx <inElementsSize){
    int4 dijnz1 = d4_IJNZ[idx0];                    
    float4 dv11 = d4_v1[idx0];
    float4 du31 = make_float4(0,0,0,0);
    float4 du32 = make_float4(0,0,0,0);

    int srcid = dijnz1.x % 10;
    if (srcid){
      du31 = d4_srcs[(srcid-1)*2];
      du32 = d4_srcs[(srcid-1)*2+1];
    }

    float4 ff1;

    //make sure all argument are the same
    ff1.x = ComputeFF(idx, dijnz1.x/10, d4_v1, inElementsSize, dv11.x, d4_AFF, Tss, du31.x, 0);
    ff1.y = ComputeFF(idx, dijnz1.y, d4_v1, inElementsSize, dv11.y, d4_AFF, Tss, du31.y, 1);
    ff1.z = ComputeFF(idx, dijnz1.z, d4_v1, inElementsSize, dv11.z, d4_AFF, Tss, du31.z, 2);
    ff1.w = ComputeFF(idx, dijnz1.w, d4_v1, inElementsSize, dv11.w, d4_AFF, Tss, du31.w, 3);
    ff4[idx0] = ff1;

    dijnz1 = d4_IJNZ[idx0+1];
    dv11   = d4_v1[idx0+1];
    ff1.x = ComputeFF(idx, dijnz1.x, d4_v1, inElementsSize, dv11.x, d4_AFF, Tss, du32.x, 4);
    ff1.y = ComputeFF(idx, dijnz1.y, d4_v1, inElementsSize, dv11.y, d4_AFF, Tss, du32.y, 5);
    ff1.z = ComputeFF(idx, dijnz1.z, d4_v1, inElementsSize, dv11.z, d4_AFF, Tss, du32.z, 6);
    ff1.w = ComputeFF(idx, dijnz1.w, d4_v1, inElementsSize, dv11.w, d4_AFF, Tss, du32.w, 7);
    ff4[idx0+1] = ff1;

  }
}

__global__ void  UpdateSolution(int inElementsSize, float4 *d4_AL, float4 *d4_v1, float4 *ff4, float4 *diff4_v1){

  //TODO: use the temproray ff4 array to hold the temprary diff4_v1 data 
  int idx = blockIdx.x * blockDim.x + threadIdx.x;

  float4 dff1 = ff4[idx*2];
  float4 dff2 = ff4[idx*2+1];
  float4 dvvv;
  int idx0 = idx*16;

  float4 diffv1 = d4_v1[idx*2];

  dvvv.x = DotProdVec4(d4_AL[idx0],   dff1) + DotProdVec4(d4_AL[idx0+1], dff2);
  dvvv.y = DotProdVec4(d4_AL[idx0+2], dff1) + DotProdVec4(d4_AL[idx0+2+1], dff2);
  dvvv.z = DotProdVec4(d4_AL[idx0+4], dff1) + DotProdVec4(d4_AL[idx0+4+1], dff2);
  dvvv.w = DotProdVec4(d4_AL[idx0+6], dff1) + DotProdVec4(d4_AL[idx0+6+1], dff2);
  d4_v1[idx*2] = dvvv;

  diffv1.x = fabs(diffv1.x - dvvv.x);
  diffv1.y = fabs(diffv1.y - dvvv.y);
  diffv1.z = fabs(diffv1.z - dvvv.z);
  diffv1.w = fabs(diffv1.w - dvvv.w);
  diff4_v1[idx*2] = diffv1;

  dvvv.x = DotProdVec4(d4_AL[idx0+8],  dff1) + DotProdVec4(d4_AL[idx0+8+1], dff2);
  dvvv.y = DotProdVec4(d4_AL[idx0+10], dff1) + DotProdVec4(d4_AL[idx0+10+1], dff2);
  dvvv.z = DotProdVec4(d4_AL[idx0+12], dff1) + DotProdVec4(d4_AL[idx0+12+1], dff2);
  dvvv.w = DotProdVec4(d4_AL[idx0+14], dff1) + DotProdVec4(d4_AL[idx0+14+1], dff2);

  diffv1 = d4_v1[idx*2+1];
  d4_v1[idx*2+1] = dvvv;

  diffv1.x = fabs(diffv1.x - dvvv.x);
  diffv1.y = fabs(diffv1.y - dvvv.y);
  diffv1.z = fabs(diffv1.z - dvvv.z);
  diffv1.w = fabs(diffv1.w - dvvv.w);
  diff4_v1[idx*2+1] = diffv1;

}

int InitializeSolver( int device){

  //TODO: move all cuda initialization here 
  
  static int setdev = 0;

  int deviceCount;
  hipGetDeviceCount(&deviceCount);
  if (deviceCount == 0){
    fprintf(stderr, "There is no device supporting CUDA ");
    return 1;
  }
  
  if (device < 0 || device >= deviceCount){
    fprintf(stderr, "No such cuda device = %d \n", device);
    return 1;
  }

  else {
    if (!setdev){
      if( hipSetDevice( device ) != hipSuccess) {
	fprintf(stderr, "Failed setting cuda device \n ");
	return 1;
      }
      setdev = 1;
    }
  }

  char host[HOST_NAME_SIZE];
  gethostname(host, HOST_NAME_SIZE);

  printf("[%s] using id=%d cuda device.\n", host, device);

  return 0; 

}

extern "C" int SolveVaiCuda(int inElementsSize, int* IJNZ, float *v1, float *AFF, float *AL,
			    int device, float Tss, float *srcs, int num_srcs, float scaledTol, 
			    int maxNumIterations, int checkMax, float eps, int printFlag){

  static int setdev = 0;
  if (setdev == 0){
    InitializeSolver(device);
    setdev = 1;
  }

  int   iterations, check = 0; 
  float loop_step_diff = FLT_MAX, diff = FLT_MAX;

  dim3 blocks(THREADS_PER_BLOCK);
  dim3 grid(inElementsSize/blocks.x + 1);

  float4 *d4_v1, *d4_AL, *d4_AFF, *d4_ff, *d4_srcs, *diff4_v1;
  int4   *d4_IJNZ;

  float * ff = (float *) malloc(inElementsSize*8*sizeof(float));

  assert( hipMalloc( (void**) &d4_IJNZ, inElementsSize*2*sizeof(int4))   == hipSuccess);
  assert( hipMemcpy( d4_IJNZ, IJNZ, inElementsSize*8*sizeof(int), hipMemcpyHostToDevice) == hipSuccess);

  assert( hipMalloc( (void**) &d4_AFF, inElementsSize*16*sizeof(float4))   == hipSuccess);
  assert( hipMemcpy( d4_AFF, AFF, inElementsSize*64*sizeof(float), hipMemcpyHostToDevice) == hipSuccess);

  assert( hipMalloc( (void**) &d4_AL, inElementsSize*16*sizeof(float4))   == hipSuccess);
  assert( hipMemcpy( d4_AL, AL, inElementsSize*64*sizeof(float), hipMemcpyHostToDevice) == hipSuccess);

  assert( hipMalloc( (void**) &d4_ff, inElementsSize*2*sizeof(float4))   == hipSuccess);

  assert( hipMalloc( (void**) &d4_srcs, num_srcs*2*sizeof(float4))   == hipSuccess);
  assert( hipMemcpy( d4_srcs, srcs, num_srcs*8*sizeof(float), hipMemcpyHostToDevice) == hipSuccess);

  assert( hipMalloc( (void**) &d4_v1, inElementsSize*2*sizeof(float4))   == hipSuccess);
  assert( hipMemset( d4_v1, 0, inElementsSize*2*sizeof(float4)) == hipSuccess);

  assert( hipMalloc( (void**) &diff4_v1, inElementsSize*2*sizeof(float4))   == hipSuccess);
  assert( hipMemset( diff4_v1, -1, inElementsSize*8*sizeof(float)) == hipSuccess);

  ///temp stuff
  assert( hipMemset( d4_ff, 0,  inElementsSize*8*sizeof(float)) == hipSuccess);
  
  for(iterations = 0; iterations < maxNumIterations && loop_step_diff > scaledTol && 
	check < checkMax; iterations++) {

    DevComputeff4<<<grid, blocks>>>(inElementsSize, d4_IJNZ, d4_srcs, d4_AFF, d4_v1, d4_ff, Tss);
    UpdateSolution<<<grid, blocks>>> (inElementsSize, d4_AL, d4_v1, d4_ff, diff4_v1);

    loop_step_diff = full_reduce((float*)diff4_v1, (float*) diff4_v1, inElementsSize*8, 512, 256);
    if (fabs(loop_step_diff - diff) < eps ) check++;
    else {
      check = 0;
      diff = loop_step_diff;
    }
  }

  assert(hipMemcpy(v1, d4_v1, (inElementsSize)*2*sizeof(float4), hipMemcpyDeviceToHost) == hipSuccess);

  /*
  cudaError_t err = cudaMemcpy(v1, d4_v1, (inElementsSize)*2*sizeof(float4), cudaMemcpyDeviceToHost);

  if (err ==     cudaSuccess)
    printf("Success ");
  else if (err == cudaErrorInvalidValue)
    printf("cudaErrorInvalidValue ");
  else if (err == cudaErrorInvalidDevicePointer)
    printf("cudaErrorInvalidDevicePointer ");
  else if (err == cudaErrorInvalidMemcpyDirection)
    printf("cudaErrorInvalidMemcpyDirection ");
  else 
    printf("Unknown error ");
  */


  hipDeviceSynchronize();

  hipFree(d4_IJNZ);
  hipFree(d4_AFF);
  hipFree(d4_AL);
  hipFree(d4_ff);
  hipFree(d4_srcs);
  hipFree(diff4_v1);
  hipFree(d4_v1);

  return iterations;
  
}

	.text
	.file	"poisson_vai_cuda.hip"
	.globl	_Z28__device_stub__DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f # -- Begin function _Z28__device_stub__DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f
	.type	_Z28__device_stub__DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f,@function
_Z28__device_stub__DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f: # @_Z28__device_stub__DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f
	.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	$160, %rsp
	.cfi_def_cfa_offset 208
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	4(%rsp), %rax
	movl	%edi, (%rax)
	leaq	56(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	48(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	40(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	32(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	24(%rsp), %r8
	movq	%r9, (%r8)
	movq	%rsp, %r9
	movss	%xmm0, (%r9)
	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	%r9, 48(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z13DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$176, %rsp
	.cfi_adjust_cfa_offset -176
	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	_Z28__device_stub__DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f, .Lfunc_end0-_Z28__device_stub__DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f
	.cfi_endproc
                                        # -- End function
	.globl	_Z29__device_stub__UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_ # -- Begin function _Z29__device_stub__UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_
	.type	_Z29__device_stub__UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_,@function
_Z29__device_stub__UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_: # @_Z29__device_stub__UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_
	.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	$144, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	12(%rsp), %rax
	movl	%edi, (%rax)
	leaq	56(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	48(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	40(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	32(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z14UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$160, %rsp
	.cfi_adjust_cfa_offset -160
	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	_Z29__device_stub__UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_, .Lfunc_end1-_Z29__device_stub__UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_
	.cfi_endproc
                                        # -- End function
	.globl	_Z16InitializeSolveri           # -- Begin function _Z16InitializeSolveri
	.type	_Z16InitializeSolveri,@function
_Z16InitializeSolveri:                  # @_Z16InitializeSolveri
	.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	$224, %rsp
	.cfi_def_cfa_offset 256
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %rbp, -16
	movl	%edi, %ebx
	leaq	12(%rsp), %r14
	movq	%r14, %rdi
	callq	hipGetDeviceCount
	movl	(%r14), %eax
	testl	%eax, %eax
	je	.LBB2_8
# %bb.1:
	testl	%ebx, %ebx
	js	.LBB2_9
# %bb.2:
	cmpl	%ebx, %eax
	jle	.LBB2_9
# %bb.3:
	cmpb	$0, _ZZ16InitializeSolveriE6setdev(%rip)
	jne	.LBB2_6
# %bb.4:
	movl	%ebx, %edi
	callq	hipSetDevice
	testl	%eax, %eax
	jne	.LBB2_10
# %bb.5:
	movb	$1, _ZZ16InitializeSolveriE6setdev(%rip)
.LBB2_6:
	leaq	16(%rsp), %r14
	movl	$200, %esi
	movq	%r14, %rdi
	callq	gethostname
	xorl	%ebp, %ebp
	movl	$.L.str.3, %edi
	movq	%r14, %rsi
	movl	%ebx, %edx
	xorl	%eax, %eax
	callq	printf
.LBB2_7:
	movl	%ebp, %eax
	addq	$224, %rsp
	.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
.LBB2_8:
	.cfi_def_cfa_offset 256
	movq	stderr(%rip), %rcx
	movl	$.L.str, %edi
	movl	$35, %esi
	jmp	.LBB2_11
.LBB2_9:
	movq	stderr(%rip), %rdi
	movl	$.L.str.1, %esi
	movl	%ebx, %edx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB2_12
.LBB2_10:
	movq	stderr(%rip), %rcx
	movl	$.L.str.2, %edi
	movl	$29, %esi
.LBB2_11:
	movl	$1, %edx
	callq	fwrite@PLT
.LBB2_12:
	movl	$1, %ebp
	jmp	.LBB2_7
.Lfunc_end2:
	.size	_Z16InitializeSolveri, .Lfunc_end2-_Z16InitializeSolveri
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function SolveVaiCuda
.LCPI3_0:
	.long	0x7f7fffff                      # float 3.40282347E+38
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0
.LCPI3_1:
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
	.text
	.globl	SolveVaiCuda
	.type	SolveVaiCuda,@function
SolveVaiCuda:                           # @SolveVaiCuda
	.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	$40, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movss	%xmm2, 8(%rsp)                  # 4-byte Spill
	movss	%xmm1, 4(%rsp)                  # 4-byte Spill
	movss	%xmm0, 12(%rsp)                 # 4-byte Spill
	movl	%edi, %r14d
	cmpb	$0, _ZZ12SolveVaiCudaE6setdev(%rip)
	jne	.LBB3_2
# %bb.1:
	movl	%r9d, %edi
	callq	_Z16InitializeSolveri
	movb	$1, _ZZ12SolveVaiCudaE6setdev(%rip)
.LBB3_2:
	xorl	%ebx, %ebx
	movss	.LCPI3_0(%rip), %xmm1           # xmm1 = mem[0],zero,zero,zero
	ucomiss	4(%rsp), %xmm1                  # 4-byte Folded Reload
	jbe	.LBB3_13
# %bb.3:
	cmpl	$0, 112(%rsp)
	jle	.LBB3_13
# %bb.4:
	cmpl	$0, 120(%rsp)
	jle	.LBB3_13
# %bb.5:                                # %.lr.ph.preheader
	movabsq	$4294967328, %r15               # imm = 0x100000020
	movl	%r14d, %eax
	shrl	$5, %eax
	leaq	(%rax,%r15), %r12
	addq	$-31, %r12
	leal	(,%r14,8), %ebp
	xorl	%ebx, %ebx
	xorl	%r13d, %r13d
.LBB3_6:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movaps	%xmm1, 16(%rsp)                 # 16-byte Spill
	movq	%r12, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_8
# %bb.7:                                #   in Loop: Header=BB3_6 Depth=1
	movl	%r14d, %edi
	movss	12(%rsp), %xmm0                 # 4-byte Reload
                                        # xmm0 = mem[0],zero,zero,zero
	callq	_Z28__device_stub__DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f
.LBB3_8:                                #   in Loop: Header=BB3_6 Depth=1
	movq	%r12, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_10
# %bb.9:                                #   in Loop: Header=BB3_6 Depth=1
	movl	%r14d, %edi
	callq	_Z29__device_stub__UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_
.LBB3_10:                               #   in Loop: Header=BB3_6 Depth=1
	movl	%ebp, %edx
	movl	$512, %ecx                      # imm = 0x200
	movl	$256, %r8d                      # imm = 0x100
	callq	_Z11full_reducePfS_jjj
	movaps	%xmm0, %xmm1
	movaps	16(%rsp), %xmm3                 # 16-byte Reload
	subss	%xmm3, %xmm1
	andps	.LCPI3_1(%rip), %xmm1
	incl	%r13d
	movss	8(%rsp), %xmm2                  # 4-byte Reload
                                        # xmm2 = mem[0],zero,zero,zero
	ucomiss	%xmm1, %xmm2
	movl	$0, %eax
	cmovbel	%eax, %r13d
	incl	%ebx
	ucomiss	4(%rsp), %xmm0                  # 4-byte Folded Reload
	jbe	.LBB3_13
# %bb.11:                               #   in Loop: Header=BB3_6 Depth=1
	cmpl	112(%rsp), %ebx
	jge	.LBB3_13
# %bb.12:                               #   in Loop: Header=BB3_6 Depth=1
	cmpltss	8(%rsp), %xmm1                  # 4-byte Folded Reload
	movaps	%xmm1, %xmm2
	andnps	%xmm0, %xmm2
	andps	%xmm3, %xmm1
	orps	%xmm2, %xmm1
	cmpl	120(%rsp), %r13d
	jl	.LBB3_6
.LBB3_13:                               # %.critedge
	callq	hipDeviceSynchronize
	callq	hipFree
	callq	hipFree
	callq	hipFree
	callq	hipFree
	callq	hipFree
	callq	hipFree
	callq	hipFree
	movl	%ebx, %eax
	addq	$40, %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_end3:
	.size	SolveVaiCuda, .Lfunc_end3-SolveVaiCuda
	.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	$_Z13DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f, %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	$_Z14UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_, %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	_Z13DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f,@object # @_Z13DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f
	.section	.rodata,"a",@progbits
	.globl	_Z13DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f
	.p2align	3, 0x0
_Z13DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f:
	.quad	_Z28__device_stub__DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f
	.size	_Z13DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f, 8

	.type	_Z14UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_,@object # @_Z14UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_
	.globl	_Z14UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_
	.p2align	3, 0x0
_Z14UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_:
	.quad	_Z29__device_stub__UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_
	.size	_Z14UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_, 8

	.type	_ZZ16InitializeSolveriE6setdev,@object # @_ZZ16InitializeSolveriE6setdev
	.local	_ZZ16InitializeSolveriE6setdev
	.comm	_ZZ16InitializeSolveriE6setdev,1,4
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"There is no device supporting CUDA "
	.size	.L.str, 36

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"No such cuda device = %d \n"
	.size	.L.str.1, 27

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"Failed setting cuda device \n "
	.size	.L.str.2, 30

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"[%s] using id=%d cuda device.\n"
	.size	.L.str.3, 31

	.type	_ZZ12SolveVaiCudaE6setdev,@object # @_ZZ12SolveVaiCudaE6setdev
	.local	_ZZ12SolveVaiCudaE6setdev
	.comm	_ZZ12SolveVaiCudaE6setdev,1,4
	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z13DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f"
	.size	.L__unnamed_1, 64

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z14UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_"
	.size	.L__unnamed_2, 54

	.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 _Z28__device_stub__DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f
	.addrsig_sym _Z29__device_stub__UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z13DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f
	.addrsig_sym _Z14UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z13DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f ; -- Begin function _Z13DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f
	.globl	_Z13DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f
	.p2align	8
	.type	_Z13DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f,@function
_Z13DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f: ; @_Z13DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x44
	s_load_b32 s12, s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[15:16], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_i32_e64 s12, v15
	s_cbranch_execz .LBB0_20
; %bb.1:
	s_load_b256 s[4:11], s[0:1], 0x8
	v_dual_mov_b32 v23, 0 :: v_dual_lshlrev_b32 v0, 1, v15
	s_load_b64 s[2:3], s[0:1], 0x28
	v_mov_b32_e32 v8, 0
	v_mov_b32_e32 v10, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_ashrrev_i32_e32 v1, 31, v0
	v_dual_mov_b32 v2, 0 :: v_dual_mov_b32 v9, 0
	v_mov_b32_e32 v3, 0
	v_lshlrev_b64 v[21:22], 4, v[0:1]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v17, vcc_lo, s4, v21
	v_add_co_ci_u32_e32 v18, vcc_lo, s5, v22, vcc_lo
	v_add_co_u32 v19, vcc_lo, s10, v21
	v_add_co_ci_u32_e32 v20, vcc_lo, s11, v22, vcc_lo
	global_load_b128 v[11:14], v[17:18], off
	s_mov_b32 s4, exec_lo
	global_load_b128 v[4:7], v[19:20], off
	s_waitcnt vmcnt(1)
	v_mul_hi_i32 v0, 0x66666667, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshrrev_b32_e32 v1, 31, v0
	v_ashrrev_i32_e32 v0, 2, v0
	v_dual_mov_b32 v1, 0 :: v_dual_add_nc_u32 v16, v0, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v0, v16, 10
	v_sub_nc_u32_e32 v24, v11, v0
	v_dual_mov_b32 v11, 0 :: v_dual_mov_b32 v0, 0
	s_delay_alu instid0(VALU_DEP_2)
	v_cmpx_ne_u32_e32 0, v24
; %bb.2:
	v_lshlrev_b32_e32 v0, 1, v24
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v1, 31, v0
	v_lshlrev_b64 v[0:1], 4, v[0:1]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s6, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s7, v1, vcc_lo
	s_clause 0x1
	global_load_b128 v[8:11], v[0:1], off offset:-32
	global_load_b128 v[0:3], v[0:1], off offset:-16
; %bb.3:
	s_or_b32 exec_lo, exec_lo, s4
	v_dual_mov_b32 v24, 0 :: v_dual_lshlrev_b32 v15, 4, v15
	s_mov_b32 s4, exec_lo
	v_cmpx_ne_u32_e64 s12, v16
	s_cbranch_execz .LBB0_5
; %bb.4:
	v_lshlrev_b32_e32 v24, 1, v16
	v_ashrrev_i32_e32 v16, 31, v15
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v25, 31, v24
	v_lshlrev_b64 v[26:27], 4, v[15:16]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[24:25], 4, v[24:25]
	v_add_co_u32 v28, vcc_lo, s8, v26
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v29, vcc_lo, s9, v27, vcc_lo
	v_add_co_u32 v36, vcc_lo, s10, v24
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_ci_u32_e32 v37, vcc_lo, s11, v25, vcc_lo
	s_clause 0x1
	global_load_b128 v[24:27], v[28:29], off offset:16
	global_load_b128 v[28:31], v[28:29], off
	s_clause 0x1
	global_load_b128 v[32:35], v[36:37], off
	global_load_b128 v[36:39], v[36:37], off offset:16
	s_waitcnt vmcnt(0)
	v_mul_f32_e32 v25, v37, v25
	v_mul_f32_e32 v16, v33, v29
	v_add_f32_e32 v4, v4, v39
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fmac_f32_e32 v25, v36, v24
	v_fmac_f32_e32 v16, v32, v28
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mul_f32 v4, 0.5, v4 :: v_dual_fmac_f32 v25, v38, v26
	v_dual_fmac_f32 v16, v34, v30 :: v_dual_fmac_f32 v25, v39, v27
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v16, v35, v31
	v_add_f32_e32 v24, v16, v25
.LBB0_5:                                ; %_Z9ComputeFFiiP15HIP_vector_typeIfLj4EEifS1_ffi.exit
	s_or_b32 exec_lo, exec_lo, s4
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s4, exec_lo
	v_cmpx_ne_u32_e64 s12, v12
	s_cbranch_execz .LBB0_7
; %bb.6:
	v_lshlrev_b32_e32 v25, 1, v12
	v_ashrrev_i32_e32 v16, 31, v15
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v26, 31, v25
	v_lshlrev_b64 v[27:28], 4, v[15:16]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[25:26], 4, v[25:26]
	v_add_co_u32 v29, vcc_lo, s8, v27
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v30, vcc_lo, s9, v28, vcc_lo
	v_add_co_u32 v37, vcc_lo, s10, v25
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_ci_u32_e32 v38, vcc_lo, s11, v26, vcc_lo
	s_clause 0x1
	global_load_b128 v[25:28], v[29:30], off offset:48
	global_load_b128 v[29:32], v[29:30], off offset:32
	s_clause 0x1
	global_load_b128 v[33:36], v[37:38], off
	global_load_b128 v[37:40], v[37:38], off offset:16
	s_waitcnt vmcnt(1)
	v_mul_f32_e32 v12, v34, v30
	s_waitcnt vmcnt(0)
	v_dual_mul_f32 v16, v38, v26 :: v_dual_add_f32 v5, v5, v39
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v12, v33, v29
	v_fmac_f32_e32 v16, v37, v25
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_dual_mul_f32 v5, 0.5, v5 :: v_dual_fmac_f32 v12, v35, v31
	v_fmac_f32_e32 v16, v39, v27
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v12, v36, v32
	v_fmac_f32_e32 v16, v40, v28
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f32_e32 v23, v12, v16
.LBB0_7:                                ; %_Z9ComputeFFiiP15HIP_vector_typeIfLj4EEifS1_ffi.exit127
	s_or_b32 exec_lo, exec_lo, s4
	v_dual_mov_b32 v25, 0 :: v_dual_mov_b32 v26, 0
	s_mov_b32 s4, exec_lo
	v_cmpx_ne_u32_e64 s12, v13
	s_cbranch_execz .LBB0_9
; %bb.8:
	v_lshlrev_b32_e32 v12, 1, v13
	v_ashrrev_i32_e32 v16, 31, v15
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v13, 31, v12
	v_lshlrev_b64 v[26:27], 4, v[15:16]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[12:13], 4, v[12:13]
	v_add_co_u32 v30, vcc_lo, s8, v26
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v31, vcc_lo, s9, v27, vcc_lo
	v_add_co_u32 v12, vcc_lo, s10, v12
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_ci_u32_e32 v13, vcc_lo, s11, v13, vcc_lo
	s_clause 0x1
	global_load_b128 v[26:29], v[30:31], off offset:80
	global_load_b128 v[30:33], v[30:31], off offset:64
	s_clause 0x1
	global_load_b128 v[34:37], v[12:13], off
	global_load_b128 v[38:41], v[12:13], off offset:16
	s_waitcnt vmcnt(0)
	v_mul_f32_e32 v13, v39, v27
	v_mul_f32_e32 v12, v35, v31
	v_add_f32_e32 v6, v6, v39
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fmac_f32_e32 v13, v38, v26
	v_fmac_f32_e32 v12, v34, v30
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mul_f32 v6, 0.5, v6 :: v_dual_fmac_f32 v13, v40, v28
	v_dual_fmac_f32 v12, v36, v32 :: v_dual_fmac_f32 v13, v41, v29
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v12, v37, v33
	v_add_f32_e32 v26, v12, v13
.LBB0_9:                                ; %_Z9ComputeFFiiP15HIP_vector_typeIfLj4EEifS1_ffi.exit151
	s_or_b32 exec_lo, exec_lo, s4
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s4, exec_lo
	v_cmpx_ne_u32_e64 s12, v14
	s_cbranch_execz .LBB0_11
; %bb.10:
	v_lshlrev_b32_e32 v12, 1, v14
	v_ashrrev_i32_e32 v16, 31, v15
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v13, 31, v12
	v_lshlrev_b64 v[27:28], 4, v[15:16]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[12:13], 4, v[12:13]
	v_add_co_u32 v31, vcc_lo, s8, v27
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v32, vcc_lo, s9, v28, vcc_lo
	v_add_co_u32 v12, vcc_lo, s10, v12
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_ci_u32_e32 v13, vcc_lo, s11, v13, vcc_lo
	s_clause 0x1
	global_load_b128 v[27:30], v[31:32], off offset:112
	global_load_b128 v[31:34], v[31:32], off offset:96
	s_clause 0x1
	global_load_b128 v[35:38], v[12:13], off
	global_load_b128 v[39:42], v[12:13], off offset:16
	s_waitcnt vmcnt(0)
	v_mul_f32_e32 v13, v40, v28
	v_dual_mul_f32 v12, v36, v32 :: v_dual_add_f32 v7, v7, v39
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v13, v39, v27
	v_fmac_f32_e32 v12, v35, v31
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_f32_e32 v7, 0.5, v7
	v_fmac_f32_e32 v13, v41, v29
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_fmac_f32 v12, v37, v33 :: v_dual_fmac_f32 v13, v42, v30
	v_fmac_f32_e32 v12, v38, v34
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f32_e32 v25, v12, v13
.LBB0_11:                               ; %_Z9ComputeFFiiP15HIP_vector_typeIfLj4EEifS1_ffi.exit175
	s_or_b32 exec_lo, exec_lo, s4
	s_load_b32 s0, s[0:1], 0x30
	v_add_co_u32 v12, vcc_lo, s2, v21
	v_add_co_ci_u32_e32 v13, vcc_lo, s3, v22, vcc_lo
	v_mov_b32_e32 v14, 0
	s_mov_b32 s1, exec_lo
	s_waitcnt vmcnt(0) lgkmcnt(0)
	v_dual_fmac_f32 v10, s0, v6 :: v_dual_fmac_f32 v9, s0, v5
	v_dual_fmac_f32 v8, s0, v4 :: v_dual_fmac_f32 v11, s0, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_dual_sub_f32 v6, v10, v26 :: v_dual_sub_f32 v5, v9, v23
	v_dual_sub_f32 v4, v8, v24 :: v_dual_sub_f32 v7, v11, v25
	global_store_b128 v[12:13], v[4:7], off
	global_load_b128 v[8:11], v[17:18], off offset:16
	global_load_b128 v[4:7], v[19:20], off offset:16
	v_mov_b32_e32 v17, 0
	s_waitcnt vmcnt(1)
	v_cmpx_ne_u32_e64 s12, v8
	s_cbranch_execz .LBB0_13
; %bb.12:
	v_lshlrev_b32_e32 v17, 1, v8
	v_ashrrev_i32_e32 v16, 31, v15
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v18, 31, v17
	v_lshlrev_b64 v[19:20], 4, v[15:16]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[16:17], 4, v[17:18]
	v_add_co_u32 v21, vcc_lo, s8, v19
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v22, vcc_lo, s9, v20, vcc_lo
	v_add_co_u32 v28, vcc_lo, s10, v16
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_ci_u32_e32 v29, vcc_lo, s11, v17, vcc_lo
	s_clause 0x1
	global_load_b128 v[16:19], v[21:22], off offset:144
	global_load_b128 v[20:23], v[21:22], off offset:128
	s_clause 0x1
	global_load_b128 v[24:27], v[28:29], off
	global_load_b128 v[28:31], v[28:29], off offset:16
	s_waitcnt vmcnt(1)
	v_mul_f32_e32 v8, v25, v21
	s_waitcnt vmcnt(0)
	v_dual_mul_f32 v17, v29, v17 :: v_dual_add_f32 v4, v4, v27
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v8, v24, v20
	v_fmac_f32_e32 v17, v28, v16
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_f32_e32 v4, 0.5, v4
	v_fmac_f32_e32 v8, v26, v22
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_fmac_f32 v17, v30, v18 :: v_dual_fmac_f32 v8, v27, v23
	v_fmac_f32_e32 v17, v31, v19
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f32_e32 v17, v8, v17
.LBB0_13:                               ; %_Z9ComputeFFiiP15HIP_vector_typeIfLj4EEifS1_ffi.exit199
	s_or_b32 exec_lo, exec_lo, s1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s1, exec_lo
	v_cmpx_ne_u32_e64 s12, v9
	s_cbranch_execz .LBB0_15
; %bb.14:
	v_lshlrev_b32_e32 v8, 1, v9
	v_ashrrev_i32_e32 v16, 31, v15
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v9, 31, v8
	v_lshlrev_b64 v[18:19], 4, v[15:16]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[8:9], 4, v[8:9]
	v_add_co_u32 v22, vcc_lo, s8, v18
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v23, vcc_lo, s9, v19, vcc_lo
	v_add_co_u32 v8, vcc_lo, s10, v8
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_ci_u32_e32 v9, vcc_lo, s11, v9, vcc_lo
	s_clause 0x1
	global_load_b128 v[18:21], v[22:23], off offset:176
	global_load_b128 v[22:25], v[22:23], off offset:160
	s_clause 0x1
	global_load_b128 v[26:29], v[8:9], off
	global_load_b128 v[30:33], v[8:9], off offset:16
	s_waitcnt vmcnt(0)
	v_mul_f32_e32 v9, v31, v19
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mul_f32 v8, v27, v23 :: v_dual_fmac_f32 v9, v30, v18
	v_dual_fmac_f32 v8, v26, v22 :: v_dual_fmac_f32 v9, v32, v20
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_fmac_f32 v8, v28, v24 :: v_dual_fmac_f32 v9, v33, v21
	v_fmac_f32_e32 v8, v29, v25
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_add_f32 v5, v5, v28 :: v_dual_add_f32 v14, v8, v9
	v_mul_f32_e32 v5, 0.5, v5
.LBB0_15:                               ; %_Z9ComputeFFiiP15HIP_vector_typeIfLj4EEifS1_ffi.exit223
	s_or_b32 exec_lo, exec_lo, s1
	v_dual_mov_b32 v8, 0 :: v_dual_mov_b32 v9, 0
	s_mov_b32 s1, exec_lo
	v_cmpx_ne_u32_e64 s12, v10
	s_cbranch_execz .LBB0_17
; %bb.16:
	v_lshlrev_b32_e32 v9, 1, v10
	v_ashrrev_i32_e32 v16, 31, v15
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v10, 31, v9
	v_lshlrev_b64 v[18:19], 4, v[15:16]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[9:10], 4, v[9:10]
	v_add_co_u32 v22, vcc_lo, s8, v18
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v23, vcc_lo, s9, v19, vcc_lo
	v_add_co_u32 v9, vcc_lo, s10, v9
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_ci_u32_e32 v10, vcc_lo, s11, v10, vcc_lo
	s_clause 0x1
	global_load_b128 v[18:21], v[22:23], off offset:208
	global_load_b128 v[22:25], v[22:23], off offset:192
	s_clause 0x1
	global_load_b128 v[26:29], v[9:10], off
	global_load_b128 v[30:33], v[9:10], off offset:16
	s_waitcnt vmcnt(1)
	v_add_f32_e32 v6, v6, v27
	s_waitcnt vmcnt(0)
	v_mul_f32_e32 v10, v31, v19
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_dual_mul_f32 v9, v27, v23 :: v_dual_mul_f32 v6, 0.5, v6
	v_fmac_f32_e32 v10, v30, v18
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_fmac_f32 v9, v26, v22 :: v_dual_fmac_f32 v10, v32, v20
	v_dual_fmac_f32 v9, v28, v24 :: v_dual_fmac_f32 v10, v33, v21
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v9, v29, v25
	v_add_f32_e32 v9, v9, v10
.LBB0_17:                               ; %_Z9ComputeFFiiP15HIP_vector_typeIfLj4EEifS1_ffi.exit247
	s_or_b32 exec_lo, exec_lo, s1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s1, exec_lo
	v_cmpx_ne_u32_e64 s12, v11
	s_cbranch_execz .LBB0_19
; %bb.18:
	v_lshlrev_b32_e32 v10, 1, v11
	v_ashrrev_i32_e32 v16, 31, v15
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v11, 31, v10
	v_lshlrev_b64 v[15:16], 4, v[15:16]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[10:11], 4, v[10:11]
	v_add_co_u32 v15, vcc_lo, s8, v15
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v16, vcc_lo, s9, v16, vcc_lo
	v_add_co_u32 v10, vcc_lo, s10, v10
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_ci_u32_e32 v11, vcc_lo, s11, v11, vcc_lo
	s_clause 0x1
	global_load_b128 v[18:21], v[15:16], off offset:240
	global_load_b128 v[22:25], v[15:16], off offset:224
	s_clause 0x1
	global_load_b128 v[26:29], v[10:11], off
	global_load_b128 v[30:33], v[10:11], off offset:16
	s_waitcnt vmcnt(1)
	v_add_f32_e32 v7, v7, v26
	v_mul_f32_e32 v8, v27, v23
	s_waitcnt vmcnt(0)
	v_mul_f32_e32 v10, v31, v19
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_dual_mul_f32 v7, 0.5, v7 :: v_dual_fmac_f32 v8, v26, v22
	v_fmac_f32_e32 v10, v30, v18
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v8, v28, v24
	v_fmac_f32_e32 v10, v32, v20
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v8, v29, v25
	v_fmac_f32_e32 v10, v33, v21
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f32_e32 v8, v8, v10
.LBB0_19:                               ; %_Z9ComputeFFiiP15HIP_vector_typeIfLj4EEifS1_ffi.exit271
	s_or_b32 exec_lo, exec_lo, s1
	s_waitcnt vmcnt(0)
	v_dual_fmac_f32 v2, s0, v6 :: v_dual_fmac_f32 v1, s0, v5
	v_dual_fmac_f32 v0, s0, v4 :: v_dual_fmac_f32 v3, s0, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_dual_sub_f32 v2, v2, v9 :: v_dual_sub_f32 v1, v1, v14
	v_dual_sub_f32 v0, v0, v17 :: v_dual_sub_f32 v3, v3, v8
	global_store_b128 v[12:13], v[0:3], off offset:16
.LBB0_20:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z13DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 312
		.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 43
		.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	_Z13DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f, .Lfunc_end0-_Z13DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 2032
; NumSgprs: 18
; NumVgprs: 43
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 5
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 43
; Occupancy: 16
; WaveLimiterHint : 1
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z14UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_ ; -- Begin function _Z14UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_
	.globl	_Z14UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_
	.p2align	8
	.type	_Z14UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_,@function
_Z14UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_: ; @_Z14UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_
; %bb.0:
	s_load_b32 s2, s[0:1], 0x34
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_load_b256 s[0:7], s[0:1], 0x8
	v_lshlrev_b32_e32 v0, 1, v1
	v_lshlrev_b32_e32 v2, 4, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v1, 31, v0
	v_ashrrev_i32_e32 v3, 31, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[44:45], 4, v[0:1]
	v_lshlrev_b64 v[0:1], 4, v[2:3]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v4, vcc_lo, s4, v44
	v_add_co_ci_u32_e32 v5, vcc_lo, s5, v45, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v46, vcc_lo, s0, v0
	v_add_co_ci_u32_e32 v47, vcc_lo, s1, v1, vcc_lo
	s_clause 0x1
	global_load_b128 v[0:3], v[4:5], off offset:16
	global_load_b128 v[4:7], v[4:5], off
	s_clause 0x7
	global_load_b128 v[8:11], v[46:47], off
	global_load_b128 v[12:15], v[46:47], off offset:16
	global_load_b128 v[16:19], v[46:47], off offset:32
	global_load_b128 v[20:23], v[46:47], off offset:48
	global_load_b128 v[24:27], v[46:47], off offset:64
	global_load_b128 v[28:31], v[46:47], off offset:80
	global_load_b128 v[32:35], v[46:47], off offset:96
	global_load_b128 v[36:39], v[46:47], off offset:112
	v_add_co_u32 v48, vcc_lo, s2, v44
	v_add_co_ci_u32_e32 v49, vcc_lo, s3, v45, vcc_lo
	v_add_co_u32 v44, vcc_lo, s6, v44
	v_add_co_ci_u32_e32 v45, vcc_lo, s7, v45, vcc_lo
	global_load_b128 v[40:43], v[48:49], off
	s_waitcnt vmcnt(8)
	v_mul_f32_e32 v9, v5, v9
	s_waitcnt vmcnt(4)
	v_mul_f32_e32 v25, v5, v25
	s_waitcnt vmcnt(3)
	v_mul_f32_e32 v29, v1, v29
	s_waitcnt vmcnt(2)
	v_mul_f32_e32 v33, v5, v33
	v_fmac_f32_e32 v9, v4, v8
	s_waitcnt vmcnt(1)
	v_mul_f32_e32 v37, v1, v37
	v_fmac_f32_e32 v25, v4, v24
	v_fmac_f32_e32 v29, v0, v28
	v_fmac_f32_e32 v9, v6, v10
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fmac_f32_e32 v37, v0, v36
	v_fmac_f32_e32 v25, v6, v26
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fmac_f32_e32 v29, v2, v30
	v_fmac_f32_e32 v9, v7, v11
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fmac_f32_e32 v37, v2, v38
	v_fmac_f32_e32 v25, v7, v27
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v29, v3, v31
	v_fmac_f32_e32 v33, v4, v32
	v_dual_fmac_f32 v37, v3, v39 :: v_dual_add_f32 v10, v25, v29
	v_mul_f32_e32 v13, v1, v13
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v33, v6, v34
	v_fmac_f32_e32 v13, v0, v12
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v33, v7, v35
	v_fmac_f32_e32 v13, v2, v14
	s_waitcnt vmcnt(0)
	v_dual_sub_f32 v14, v42, v10 :: v_dual_mul_f32 v21, v1, v21
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f32_e32 v11, v33, v37
	v_dual_fmac_f32 v13, v3, v15 :: v_dual_and_b32 v14, 0x7fffffff, v14
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v21, v0, v20
	v_dual_sub_f32 v15, v43, v11 :: v_dual_add_f32 v8, v9, v13
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v21, v2, v22
	v_dual_sub_f32 v12, v40, v8 :: v_dual_and_b32 v15, 0x7fffffff, v15
	v_mul_f32_e32 v17, v5, v17
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_dual_fmac_f32 v21, v3, v23 :: v_dual_and_b32 v12, 0x7fffffff, v12
	v_fmac_f32_e32 v17, v4, v16
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v17, v6, v18
	v_fmac_f32_e32 v17, v7, v19
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v9, v17, v21
	v_sub_f32_e32 v13, v41, v9
	s_delay_alu instid0(VALU_DEP_1)
	v_and_b32_e32 v13, 0x7fffffff, v13
	global_store_b128 v[48:49], v[8:11], off
	global_store_b128 v[44:45], v[12:15], off
	s_clause 0x7
	global_load_b128 v[8:11], v[46:47], off offset:128
	global_load_b128 v[12:15], v[46:47], off offset:144
	global_load_b128 v[16:19], v[46:47], off offset:160
	global_load_b128 v[20:23], v[46:47], off offset:176
	global_load_b128 v[24:27], v[46:47], off offset:192
	global_load_b128 v[28:31], v[46:47], off offset:208
	global_load_b128 v[32:35], v[46:47], off offset:224
	global_load_b128 v[36:39], v[46:47], off offset:240
	global_load_b128 v[40:43], v[48:49], off offset:16
	s_waitcnt vmcnt(8)
	v_mul_f32_e32 v9, v5, v9
	s_waitcnt vmcnt(7)
	v_mul_f32_e32 v13, v1, v13
	s_waitcnt vmcnt(6)
	v_mul_f32_e32 v17, v5, v17
	s_waitcnt vmcnt(5)
	v_mul_f32_e32 v21, v1, v21
	s_waitcnt vmcnt(4)
	v_mul_f32_e32 v25, v5, v25
	s_waitcnt vmcnt(3)
	v_mul_f32_e32 v29, v1, v29
	s_waitcnt vmcnt(2)
	v_mul_f32_e32 v5, v5, v33
	s_waitcnt vmcnt(1)
	v_mul_f32_e32 v33, v1, v37
	v_fmac_f32_e32 v9, v4, v8
	v_fmac_f32_e32 v13, v0, v12
	v_fmac_f32_e32 v17, v4, v16
	v_fmac_f32_e32 v21, v0, v20
	v_fmac_f32_e32 v25, v4, v24
	v_fmac_f32_e32 v29, v0, v28
	v_fmac_f32_e32 v5, v4, v32
	v_fmac_f32_e32 v33, v0, v36
	v_fmac_f32_e32 v9, v6, v10
	v_fmac_f32_e32 v13, v2, v14
	v_fmac_f32_e32 v17, v6, v18
	v_fmac_f32_e32 v21, v2, v22
	v_fmac_f32_e32 v25, v6, v26
	v_fmac_f32_e32 v29, v2, v30
	v_fmac_f32_e32 v5, v6, v34
	v_fmac_f32_e32 v33, v2, v38
	v_fmac_f32_e32 v9, v7, v11
	v_fmac_f32_e32 v13, v3, v15
	v_fmac_f32_e32 v17, v7, v19
	v_fmac_f32_e32 v21, v3, v23
	v_fmac_f32_e32 v25, v7, v27
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_3) | instid1(VALU_DEP_4)
	v_dual_fmac_f32 v29, v3, v31 :: v_dual_add_f32 v0, v9, v13
	v_fmac_f32_e32 v5, v7, v35
	v_fmac_f32_e32 v33, v3, v39
	v_add_f32_e32 v1, v17, v21
	v_add_f32_e32 v2, v25, v29
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_dual_sub_f32 v4, v40, v0 :: v_dual_add_f32 v3, v5, v33
	v_dual_sub_f32 v5, v41, v1 :: v_dual_sub_f32 v6, v42, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_dual_sub_f32 v7, v43, v3 :: v_dual_and_b32 v4, 0x7fffffff, v4
	v_and_b32_e32 v5, 0x7fffffff, v5
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_and_b32_e32 v6, 0x7fffffff, v6
	v_and_b32_e32 v7, 0x7fffffff, v7
	global_store_b128 v[48:49], v[0:3], off offset:16
	global_store_b128 v[44:45], v[4:7], off offset:16
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z14UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 296
		.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 50
		.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	_Z14UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_, .Lfunc_end1-_Z14UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 868
; NumSgprs: 18
; NumVgprs: 50
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 6
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 50
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         40
        .size:           8
        .value_kind:     global_buffer
      - .offset:         48
        .size:           4
        .value_kind:     by_value
      - .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:           _Z13DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z13DevComputeff4iP15HIP_vector_typeIiLj4EEPS_IfLj4EES3_S3_S3_f.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     43
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         44
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         52
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         54
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         56
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         58
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         60
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         62
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         104
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 296
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z14UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z14UpdateSolutioniP15HIP_vector_typeIfLj4EES1_S1_S1_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     50
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
