#include <stdio.h>
#include "cuda.h"
#define max(x,y)  ((x) > (y)? (x) : (y))
#define min(x,y)  ((x) < (y)? (x) : (y))
#define ceil(a,b) ((a) % (b) == 0 ? (a) / (b) : ((a) / (b)) + 1)

void check_error (const char* message) {
  cudaError_t error = cudaGetLastError ();
  if (error != cudaSuccess) {
    printf ("CUDA error : %s, %s\n", message, cudaGetErrorString (error));
    exit(-1);
  }
}

__global__ void hypterm (double * flux_0, double * flux_1, double * flux_2, double * flux_3, double * flux_4, double * cons_1, double * cons_2, double * cons_3, double * cons_4, double * q_1, double * q_2, double * q_3, double * q_4, double dxinv0, double dxinv1, double dxinv2, int L, int M, int N) {
  //Determining the block's indices
  int blockdim_i= (int)(blockDim.x);
  int i0 = (int)(blockIdx.x)*(blockdim_i-8);
  int i = i0 + (int)(threadIdx.x);
  int blockdim_j= (int)(blockDim.y);
  int j0 = (int)(blockIdx.y)*(blockdim_j)+4;
  int j = j0 + (int)(threadIdx.y);

  //Declarations
  double reg_cons_1_m4=0, reg_cons_1_m3=0, reg_cons_1_m2=0, reg_cons_1_m1=0, __shared__ sh_cons_1_c0[24][16], reg_cons_1_p1=0, reg_cons_1_p2=0, reg_cons_1_p3=0, reg_cons_1_p4=0;
  double reg_cons_2_m4=0, reg_cons_2_m3=0, reg_cons_2_m2=0, reg_cons_2_m1=0, __shared__ sh_cons_2_c0[24][16], reg_cons_2_p1=0, reg_cons_2_p2=0, reg_cons_2_p3=0, reg_cons_2_p4=0;
  double reg_cons_3_m4=0, reg_cons_3_m3=0, reg_cons_3_m2=0, reg_cons_3_m1=0, __shared__ sh_cons_3_c0[24][16], reg_cons_3_p1=0, reg_cons_3_p2=0, reg_cons_3_p3=0, reg_cons_3_p4=0;
  double reg_cons_4_m4=0, reg_cons_4_m3=0, reg_cons_4_m2=0, reg_cons_4_m1=0, __shared__ sh_cons_4_c0[24][16], reg_cons_4_p1=0, reg_cons_4_p2=0, reg_cons_4_p3=0, reg_cons_4_p4=0;
  double __shared__ sh_q_1_c0[24][16];
  double __shared__ sh_q_2_c0[24][16];
  double reg_q_3_m4=0, reg_q_3_m3=0, reg_q_3_m2=0, reg_q_3_m1=0, reg_q_3_c0=0, reg_q_3_p1=0, reg_q_3_p2=0, reg_q_3_p3=0, reg_q_3_p4=0;
  double reg_q_4_m4=0, reg_q_4_m3=0, reg_q_4_m2=0, reg_q_4_m1=0, __shared__ sh_q_4_c0[24][16], reg_q_4_p1=0, reg_q_4_p2=0, reg_q_4_p3=0, reg_q_4_p4=0;

  //Value Initialization
  if (j <= M-1 & i <= min (i0+blockdim_i-1, N-1)) {
    reg_cons_1_m4 = cons_1[0 + j*N + i];
    reg_cons_1_m3 = cons_1[1*M*N + j*N + i];
    reg_cons_1_m2 = cons_1[2*M*N + j*N + i];
    reg_cons_1_m1 = cons_1[3*M*N + j*N + i];
    sh_cons_1_c0[j-j0+4][i-i0] = cons_1[4*M*N + j*N + i];
    reg_cons_1_p1 = cons_1[5*M*N + j*N + i];
    reg_cons_1_p2 = cons_1[6*M*N + j*N + i];
    reg_cons_1_p3 = cons_1[7*M*N + j*N + i];
    reg_cons_2_m4 = cons_2[0 + j*N + i];
    reg_cons_2_m3 = cons_2[1*M*N + j*N + i];
    reg_cons_2_m2 = cons_2[2*M*N + j*N + i];
    reg_cons_2_m1 = cons_2[3*M*N + j*N + i];
    sh_cons_2_c0[j-j0+4][i-i0] = cons_2[4*M*N + j*N + i];
    reg_cons_2_p1 = cons_2[5*M*N + j*N + i];
    reg_cons_2_p2 = cons_2[6*M*N + j*N + i];
    reg_cons_2_p3 = cons_2[7*M*N + j*N + i];
    reg_cons_3_m4 = cons_3[0 + j*N + i];
    reg_cons_3_m3 = cons_3[1*M*N + j*N + i];
    reg_cons_3_m2 = cons_3[2*M*N + j*N + i];
    reg_cons_3_m1 = cons_3[3*M*N + j*N + i];
    sh_cons_3_c0[j-j0+4][i-i0] = cons_3[4*M*N + j*N + i];
    reg_cons_3_p1 = cons_3[5*M*N + j*N + i];
    reg_cons_3_p2 = cons_3[6*M*N + j*N + i];
    reg_cons_3_p3 = cons_3[7*M*N + j*N + i];
    reg_cons_4_m4 = cons_4[0 + j*N + i];
    reg_cons_4_m3 = cons_4[1*M*N + j*N + i];
    reg_cons_4_m2 = cons_4[2*M*N + j*N + i];
    reg_cons_4_m1 = cons_4[3*M*N + j*N + i];
    sh_cons_4_c0[j-j0+4][i-i0] = cons_4[4*M*N + j*N + i];
    reg_cons_4_p1 = cons_4[5*M*N + j*N + i];
    reg_cons_4_p2 = cons_4[6*M*N + j*N + i];
    reg_cons_4_p3 = cons_4[7*M*N + j*N + i];
    reg_q_3_m4 = q_3[0 + j*N + i];
    reg_q_3_m3 = q_3[1*M*N + j*N + i];
    reg_q_3_m2 = q_3[2*M*N + j*N + i];
    reg_q_3_m1 = q_3[3*M*N + j*N + i];
    reg_q_3_c0 = q_3[4*M*N + j*N + i];
    reg_q_3_p1 = q_3[5*M*N + j*N + i];
    reg_q_3_p2 = q_3[6*M*N + j*N + i];
    reg_q_3_p3 = q_3[7*M*N + j*N + i];
    reg_q_4_m4 = q_4[0 + j*N + i];
    reg_q_4_m3 = q_4[1*M*N + j*N + i];
    reg_q_4_m2 = q_4[2*M*N + j*N + i];
    reg_q_4_m1 = q_4[3*M*N + j*N + i];
    sh_q_4_c0[j-j0+4][i-i0] = q_4[4*M*N + j*N + i];
    reg_q_4_p1 = q_4[5*M*N + j*N + i];
    reg_q_4_p2 = q_4[6*M*N + j*N + i];
    reg_q_4_p3 = q_4[7*M*N + j*N + i];
  }
  if(threadIdx.y < 4) {
   int jj = (j-4);
   if (jj <= M-1 & i <= min (i0+blockdim_i-1, N-1)) {
    sh_cons_1_c0[jj-j0+4][i-i0] = cons_1[4*M*N + jj*N + i];
    sh_cons_2_c0[jj-j0+4][i-i0] = cons_2[4*M*N + jj*N + i];
    sh_cons_3_c0[jj-j0+4][i-i0] = cons_3[4*M*N + jj*N + i];
    sh_cons_4_c0[jj-j0+4][i-i0] = cons_4[4*M*N + jj*N + i];
    sh_q_4_c0[jj-j0+4][i-i0] = q_4[4*M*N + jj*N + i];
   }
  } else if(threadIdx.y < 8) {
   int jj = (j-4)+16;
   if (jj <= M-1 & i <= min (i0+blockdim_i-1, N-1)) {
    sh_cons_1_c0[jj-j0+4][i-i0] = cons_1[4*M*N + jj*N + i];
    sh_cons_2_c0[jj-j0+4][i-i0] = cons_2[4*M*N + jj*N + i];
    sh_cons_3_c0[jj-j0+4][i-i0] = cons_3[4*M*N + jj*N + i];
    sh_cons_4_c0[jj-j0+4][i-i0] = cons_4[4*M*N + jj*N + i];
    sh_q_4_c0[jj-j0+4][i-i0] = q_4[4*M*N + jj*N + i];
   }
  }

  //Rest of the computation
  for (int k=4; k<=L-5; ++k) {
    //Fetch new plane
    if (j <= M-1 & i <= min (i0+blockdim_i-1, N-1)) {
        reg_cons_1_p4 = cons_1[(k+4)*M*N + j*N + i];
        reg_cons_2_p4 = cons_2[(k+4)*M*N + j*N + i];
        reg_cons_3_p4 = cons_3[(k+4)*M*N + j*N + i];
        reg_cons_4_p4 = cons_4[(k+4)*M*N + j*N + i];
        sh_q_1_c0[j-j0+4][i-i0] = q_1[k*M*N + j*N + i];
        sh_q_2_c0[j-j0+4][i-i0] = q_2[k*M*N + j*N + i];
        reg_q_3_p4 = q_3[(k+4)*M*N + j*N + i];
        reg_q_4_p4 = q_4[(k+4)*M*N + j*N + i];
    }
    if(threadIdx.y < 4) {
     int jj = (j-4);
     if (jj <= M-1 & i <= min (i0+blockdim_i-1, N-1)) {
        sh_cons_1_c0[jj-j0+4][i-i0] = cons_1[k*M*N + jj*N + i];
        sh_cons_2_c0[jj-j0+4][i-i0] = cons_2[k*M*N + jj*N + i];
        sh_cons_3_c0[jj-j0+4][i-i0] = cons_3[k*M*N + jj*N + i];
        sh_cons_4_c0[jj-j0+4][i-i0] = cons_4[k*M*N + jj*N + i];
     }
    } else if(threadIdx.y < 8) {
     int jj = (j-4)+16;
     if (jj <= M-1 & i <= min (i0+blockdim_i-1, N-1)) {
        sh_cons_1_c0[jj-j0+4][i-i0] = cons_1[k*M*N + jj*N + i];
        sh_cons_2_c0[jj-j0+4][i-i0] = cons_2[k*M*N + jj*N + i];
        sh_cons_3_c0[jj-j0+4][i-i0] = cons_3[k*M*N + jj*N + i];
        sh_cons_4_c0[jj-j0+4][i-i0] = cons_4[k*M*N + jj*N + i];
     }
    } else if(threadIdx.y < 12) {
     int jj = (j-12);
     if (jj <= M-1 & i <= min (i0+blockdim_i-1, N-1)) {
        sh_q_1_c0[jj-j0+4][i-i0] = q_1[k*M*N + jj*N + i];
        sh_q_2_c0[jj-j0+4][i-i0] = q_2[k*M*N + jj*N + i];
        sh_q_4_c0[jj-j0+4][i-i0] = q_4[k*M*N + jj*N + i];
     }
    } else {
     int jj = (j-12)+16;
     if (jj <= M-1 & i <= min (i0+blockdim_i-1, N-1)) {
        sh_q_1_c0[jj-j0+4][i-i0] = q_1[k*M*N + jj*N + i];
        sh_q_2_c0[jj-j0+4][i-i0] = q_2[k*M*N + jj*N + i];
        sh_q_4_c0[jj-j0+4][i-i0] = q_4[k*M*N + jj*N + i];
     }
    }
    __syncthreads ();
	double r0,r1,r2,r3,r4;
    if ((j <= M-5) & i >= max (i0+4, 4) & i <= min (i0+blockdim_i-5, N-5)) {
	r0=flux_0[k*M*N + j*N + i], r1=flux_1[k*M*N + j*N + i], r2=flux_2[k*M*N + j*N + i], r3=flux_3[k*M*N + j*N + i], r4 = flux_4[k*M*N + j*N + i];
//	double r0=0.0f,r1=0.0f,r2=0.0f,r3=0.0f,r4=0.0f;
        r0 -= (((((0.8f * (sh_cons_1_c0[j-j0+4][i-i0+1] - sh_cons_1_c0[j-j0+4][i-i0-1])) - (0.2f * (sh_cons_1_c0[j-j0+4][i-i0+2] - sh_cons_1_c0[j-j0+4][i-i0-2]))) + (0.038f * (sh_cons_1_c0[j-j0+4][i-i0+3] - sh_cons_1_c0[j-j0+4][i-i0-3]))) - (0.0035f * (sh_cons_1_c0[j-j0+4][i-i0+4] - sh_cons_1_c0[j-j0+4][i-i0-4]))) * dxinv0);
        r0 -= (((((0.8f * (sh_cons_2_c0[j-j0+4+1][i-i0] - sh_cons_2_c0[j-j0+4-1][i-i0])) - (0.2f * (sh_cons_2_c0[j-j0+4+2][i-i0] - sh_cons_2_c0[j-j0+4-2][i-i0]))) + (0.038f * (sh_cons_2_c0[j-j0+4+3][i-i0] - sh_cons_2_c0[j-j0+4-3][i-i0]))) - (0.0035f * (sh_cons_2_c0[j-j0+4+4][i-i0] - sh_cons_2_c0[j-j0+4-4][i-i0]))) * dxinv1);
        r0 -= (((((0.8f * (reg_cons_3_p1 - reg_cons_3_m1)) - (0.2f * (reg_cons_3_p2 - reg_cons_3_m2))) + (0.038f * (reg_cons_3_p3 - reg_cons_3_m3))) - (0.0035f * (reg_cons_3_p4 - reg_cons_3_m4))) * dxinv2);
	flux_0[k*M*N + j*N + i] = r0;
        r1 -= (((((0.8f * (((sh_cons_1_c0[j-j0+4][i-i0+1] * sh_q_1_c0[j-j0+4][i-i0+1]) - (sh_cons_1_c0[j-j0+4][i-i0-1] * sh_q_1_c0[j-j0+4][i-i0-1])) + (sh_q_4_c0[j-j0+4][i-i0+1] - sh_q_4_c0[j-j0+4][i-i0-1]))) - (0.2f * (((sh_cons_1_c0[j-j0+4][i-i0+2] * sh_q_1_c0[j-j0+4][i-i0+2]) - (sh_cons_1_c0[j-j0+4][i-i0-2] * sh_q_1_c0[j-j0+4][i-i0-2])) + (sh_q_4_c0[j-j0+4][i-i0+2] - sh_q_4_c0[j-j0+4][i-i0-2])))) + (0.038f * (((sh_cons_1_c0[j-j0+4][i-i0+3] * sh_q_1_c0[j-j0+4][i-i0+3]) - (sh_cons_1_c0[j-j0+4][i-i0-3] * sh_q_1_c0[j-j0+4][i-i0-3])) + (sh_q_4_c0[j-j0+4][i-i0+3] - sh_q_4_c0[j-j0+4][i-i0-3])))) - (0.0035f * (((sh_cons_1_c0[j-j0+4][i-i0+4] * sh_q_1_c0[j-j0+4][i-i0+4]) - (sh_cons_1_c0[j-j0+4][i-i0-4] * sh_q_1_c0[j-j0+4][i-i0-4])) + (sh_q_4_c0[j-j0+4][i-i0+4] - sh_q_4_c0[j-j0+4][i-i0-4])))) * dxinv0);
        r1 -= (((((0.8f * ((sh_cons_1_c0[j-j0+4+1][i-i0] * sh_q_2_c0[j-j0+4+1][i-i0]) - (sh_cons_1_c0[j-j0+4-1][i-i0] * sh_q_2_c0[j-j0+4-1][i-i0]))) - (0.2f * ((sh_cons_1_c0[j-j0+4+2][i-i0] * sh_q_2_c0[j-j0+4+2][i-i0]) - (sh_cons_1_c0[j-j0+4-2][i-i0] * sh_q_2_c0[j-j0+4-2][i-i0])))) + (0.038f * ((sh_cons_1_c0[j-j0+4+3][i-i0] * sh_q_2_c0[j-j0+4+3][i-i0]) - (sh_cons_1_c0[j-j0+4-3][i-i0] * sh_q_2_c0[j-j0+4-3][i-i0])))) - (0.0035f * ((sh_cons_1_c0[j-j0+4+4][i-i0] * sh_q_2_c0[j-j0+4+4][i-i0]) - (sh_cons_1_c0[j-j0+4-4][i-i0] * sh_q_2_c0[j-j0+4-4][i-i0])))) * dxinv1);
        r1 -= (((((0.8f * ((reg_cons_1_p1 * reg_q_3_p1) - (reg_cons_1_m1 * reg_q_3_m1))) - (0.2f * ((reg_cons_1_p2 * reg_q_3_p2) - (reg_cons_1_m2 * reg_q_3_m2)))) + (0.038f * ((reg_cons_1_p3 * reg_q_3_p3) - (reg_cons_1_m3 * reg_q_3_m3)))) - (0.0035f * ((reg_cons_1_p4 * reg_q_3_p4) - (reg_cons_1_m4 * reg_q_3_m4)))) * dxinv2);
	flux_1[k*M*N + j*N + i] = r1;
        r2 -= (((((0.8f * ((sh_cons_2_c0[j-j0+4][i-i0+1] * sh_q_1_c0[j-j0+4][i-i0+1]) - (sh_cons_2_c0[j-j0+4][i-i0-1] * sh_q_1_c0[j-j0+4][i-i0-1]))) - (0.2f * ((sh_cons_2_c0[j-j0+4][i-i0+2] * sh_q_1_c0[j-j0+4][i-i0+2]) - (sh_cons_2_c0[j-j0+4][i-i0-2] * sh_q_1_c0[j-j0+4][i-i0-2])))) + (0.038f * ((sh_cons_2_c0[j-j0+4][i-i0+3] * sh_q_1_c0[j-j0+4][i-i0+3]) - (sh_cons_2_c0[j-j0+4][i-i0-3] * sh_q_1_c0[j-j0+4][i-i0-3])))) - (0.0035f * ((sh_cons_2_c0[j-j0+4][i-i0+4] * sh_q_1_c0[j-j0+4][i-i0+4]) - (sh_cons_2_c0[j-j0+4][i-i0-4] * sh_q_1_c0[j-j0+4][i-i0-4])))) * dxinv0);
        r2 -= (((((0.8f * (((sh_cons_2_c0[j-j0+4+1][i-i0] * sh_q_2_c0[j-j0+4+1][i-i0]) - (sh_cons_2_c0[j-j0+4-1][i-i0] * sh_q_2_c0[j-j0+4-1][i-i0])) + (sh_q_4_c0[j-j0+4+1][i-i0] - sh_q_4_c0[j-j0+4-1][i-i0]))) - (0.2f * (((sh_cons_2_c0[j-j0+4+2][i-i0] * sh_q_2_c0[j-j0+4+2][i-i0]) - (sh_cons_2_c0[j-j0+4-2][i-i0] * sh_q_2_c0[j-j0+4-2][i-i0])) + (sh_q_4_c0[j-j0+4+2][i-i0] - sh_q_4_c0[j-j0+4-2][i-i0])))) + (0.038f * (((sh_cons_2_c0[j-j0+4+3][i-i0] * sh_q_2_c0[j-j0+4+3][i-i0]) - (sh_cons_2_c0[j-j0+4-3][i-i0] * sh_q_2_c0[j-j0+4-3][i-i0])) + (sh_q_4_c0[j-j0+4+3][i-i0] - sh_q_4_c0[j-j0+4-3][i-i0])))) - (0.0035f * (((sh_cons_2_c0[j-j0+4+4][i-i0] * sh_q_2_c0[j-j0+4+4][i-i0]) - (sh_cons_2_c0[j-j0+4-4][i-i0] * sh_q_2_c0[j-j0+4-4][i-i0])) + (sh_q_4_c0[j-j0+4+4][i-i0] - sh_q_4_c0[j-j0+4-4][i-i0])))) * dxinv1);
        r2 -= (((((0.8f * ((reg_cons_2_p1 * reg_q_3_p1) - (reg_cons_2_m1 * reg_q_3_m1))) - (0.2f * ((reg_cons_2_p2 * reg_q_3_p2) - (reg_cons_2_m2 * reg_q_3_m2)))) + (0.038f * ((reg_cons_2_p3 * reg_q_3_p3) - (reg_cons_2_m3 * reg_q_3_m3)))) - (0.0035f * ((reg_cons_2_p4 * reg_q_3_p4) - (reg_cons_2_m4 * reg_q_3_m4)))) * dxinv2);
	flux_2[k*M*N + j*N + i] = r2;
        r3 -= (((((0.8f * ((sh_cons_3_c0[j-j0+4][i-i0+1] * sh_q_1_c0[j-j0+4][i-i0+1]) - (sh_cons_3_c0[j-j0+4][i-i0-1] * sh_q_1_c0[j-j0+4][i-i0-1]))) - (0.2f * ((sh_cons_3_c0[j-j0+4][i-i0+2] * sh_q_1_c0[j-j0+4][i-i0+2]) - (sh_cons_3_c0[j-j0+4][i-i0-2] * sh_q_1_c0[j-j0+4][i-i0-2])))) + (0.038f * ((sh_cons_3_c0[j-j0+4][i-i0+3] * sh_q_1_c0[j-j0+4][i-i0+3]) - (sh_cons_3_c0[j-j0+4][i-i0-3] * sh_q_1_c0[j-j0+4][i-i0-3])))) - (0.0035f * ((sh_cons_3_c0[j-j0+4][i-i0+4] * sh_q_1_c0[j-j0+4][i-i0+4]) - (sh_cons_3_c0[j-j0+4][i-i0-4] * sh_q_1_c0[j-j0+4][i-i0-4])))) * dxinv0);
        r3 -= (((((0.8f * ((sh_cons_3_c0[j-j0+4+1][i-i0] * sh_q_2_c0[j-j0+4+1][i-i0]) - (sh_cons_3_c0[j-j0+4-1][i-i0] * sh_q_2_c0[j-j0+4-1][i-i0]))) - (0.2f * ((sh_cons_3_c0[j-j0+4+2][i-i0] * sh_q_2_c0[j-j0+4+2][i-i0]) - (sh_cons_3_c0[j-j0+4-2][i-i0] * sh_q_2_c0[j-j0+4-2][i-i0])))) + (0.038f * ((sh_cons_3_c0[j-j0+4+3][i-i0] * sh_q_2_c0[j-j0+4+3][i-i0]) - (sh_cons_3_c0[j-j0+4-3][i-i0] * sh_q_2_c0[j-j0+4-3][i-i0])))) - (0.0035f * ((sh_cons_3_c0[j-j0+4+4][i-i0] * sh_q_2_c0[j-j0+4+4][i-i0]) - (sh_cons_3_c0[j-j0+4-4][i-i0] * sh_q_2_c0[j-j0+4-4][i-i0])))) * dxinv1);
        r3 -= (((((0.8f * (((reg_cons_3_p1 * reg_q_3_p1) - (reg_cons_3_m1 * reg_q_3_m1)) + (reg_q_4_p1 - reg_q_4_m1))) - (0.2f * (((reg_cons_3_p2 * reg_q_3_p2) - (reg_cons_3_m2 * reg_q_3_m2)) + (reg_q_4_p2 - reg_q_4_m2)))) + (0.038f * (((reg_cons_3_p3 * reg_q_3_p3) - (reg_cons_3_m3 * reg_q_3_m3)) + (reg_q_4_p3 - reg_q_4_m3)))) - (0.0035f * (((reg_cons_3_p4 * reg_q_3_p4) - (reg_cons_3_m4 * reg_q_3_m4)) + (reg_q_4_p4 - reg_q_4_m4)))) * dxinv2);
	flux_3[k*M*N + j*N + i] = r3;
        r4 -= (((((0.8f * (((sh_cons_4_c0[j-j0+4][i-i0+1] * sh_q_1_c0[j-j0+4][i-i0+1]) - (sh_cons_4_c0[j-j0+4][i-i0-1] * sh_q_1_c0[j-j0+4][i-i0-1])) + ((sh_q_4_c0[j-j0+4][i-i0+1] * sh_q_1_c0[j-j0+4][i-i0+1]) - (sh_q_4_c0[j-j0+4][i-i0-1] * sh_q_1_c0[j-j0+4][i-i0-1])))) - (0.2f * (((sh_cons_4_c0[j-j0+4][i-i0+2] * sh_q_1_c0[j-j0+4][i-i0+2]) - (sh_cons_4_c0[j-j0+4][i-i0-2] * sh_q_1_c0[j-j0+4][i-i0-2])) + ((sh_q_4_c0[j-j0+4][i-i0+2] * sh_q_1_c0[j-j0+4][i-i0+2]) - (sh_q_4_c0[j-j0+4][i-i0-2] * sh_q_1_c0[j-j0+4][i-i0-2]))))) + (0.038f * (((sh_cons_4_c0[j-j0+4][i-i0+3] * sh_q_1_c0[j-j0+4][i-i0+3]) - (sh_cons_4_c0[j-j0+4][i-i0-3] * sh_q_1_c0[j-j0+4][i-i0-3])) + ((sh_q_4_c0[j-j0+4][i-i0+3] * sh_q_1_c0[j-j0+4][i-i0+3]) - (sh_q_4_c0[j-j0+4][i-i0-3] * sh_q_1_c0[j-j0+4][i-i0-3]))))) - (0.0035f * (((sh_cons_4_c0[j-j0+4][i-i0+4] * sh_q_1_c0[j-j0+4][i-i0+4]) - (sh_cons_4_c0[j-j0+4][i-i0-4] * sh_q_1_c0[j-j0+4][i-i0-4])) + ((sh_q_4_c0[j-j0+4][i-i0+4] * sh_q_1_c0[j-j0+4][i-i0+4]) - (sh_q_4_c0[j-j0+4][i-i0-4] * sh_q_1_c0[j-j0+4][i-i0-4]))))) * dxinv0);
        r4 -= (((((0.8f * (((sh_cons_4_c0[j-j0+4+1][i-i0] * sh_q_2_c0[j-j0+4+1][i-i0]) - (sh_cons_4_c0[j-j0+4-1][i-i0] * sh_q_2_c0[j-j0+4-1][i-i0])) + ((sh_q_4_c0[j-j0+4+1][i-i0] * sh_q_2_c0[j-j0+4+1][i-i0]) - (sh_q_4_c0[j-j0+4-1][i-i0] * sh_q_2_c0[j-j0+4-1][i-i0])))) - (0.2f * (((sh_cons_4_c0[j-j0+4+2][i-i0] * sh_q_2_c0[j-j0+4+2][i-i0]) - (sh_cons_4_c0[j-j0+4-2][i-i0] * sh_q_2_c0[j-j0+4-2][i-i0])) + ((sh_q_4_c0[j-j0+4+2][i-i0] * sh_q_2_c0[j-j0+4+2][i-i0]) - (sh_q_4_c0[j-j0+4-2][i-i0] * sh_q_2_c0[j-j0+4-2][i-i0]))))) + (0.038f * (((sh_cons_4_c0[j-j0+4+3][i-i0] * sh_q_2_c0[j-j0+4+3][i-i0]) - (sh_cons_4_c0[j-j0+4-3][i-i0] * sh_q_2_c0[j-j0+4-3][i-i0])) + ((sh_q_4_c0[j-j0+4+3][i-i0] * sh_q_2_c0[j-j0+4+3][i-i0]) - (sh_q_4_c0[j-j0+4-3][i-i0] * sh_q_2_c0[j-j0+4-3][i-i0]))))) - (0.0035f * (((sh_cons_4_c0[j-j0+4+4][i-i0] * sh_q_2_c0[j-j0+4+4][i-i0]) - (sh_cons_4_c0[j-j0+4-4][i-i0] * sh_q_2_c0[j-j0+4-4][i-i0])) + ((sh_q_4_c0[j-j0+4+4][i-i0] * sh_q_2_c0[j-j0+4+4][i-i0]) - (sh_q_4_c0[j-j0+4-4][i-i0] * sh_q_2_c0[j-j0+4-4][i-i0]))))) * dxinv1);
        r4 -= (((((0.8f * (((reg_cons_4_p1 * reg_q_3_p1) - (reg_cons_4_m1 * reg_q_3_m1)) + ((reg_q_4_p1 * reg_q_3_p1) - (reg_q_4_m1 * reg_q_3_m1)))) - (0.2f * (((reg_cons_4_p2 * reg_q_3_p2) - (reg_cons_4_m2 * reg_q_3_m2)) + ((reg_q_4_p2 * reg_q_3_p2) - (reg_q_4_m2 * reg_q_3_m2))))) + (0.038f * (((reg_cons_4_p3 * reg_q_3_p3) - (reg_cons_4_m3 * reg_q_3_m3)) + ((reg_q_4_p3 * reg_q_3_p3) - (reg_q_4_m3 * reg_q_3_m3))))) - (0.0035f * (((reg_cons_4_p4 * reg_q_3_p4) - (reg_cons_4_m4 * reg_q_3_m4)) + ((reg_q_4_p4 * reg_q_3_p4) - (reg_q_4_m4 * reg_q_3_m4))))) * dxinv2);
	flux_4[k*M*N + j*N + i] = r4;
    }
    __syncthreads ();
    //Value rotation


  if (j <= M-1 & i <= min (i0+blockdim_i-1, N-1)) {
    reg_cons_1_m4 = reg_cons_1_m3;
    reg_cons_1_m3 = reg_cons_1_m2;
    reg_cons_1_m2 = reg_cons_1_m1;
    reg_cons_1_m1 = sh_cons_1_c0[j-j0+4][i-i0];
    sh_cons_1_c0[j-j0+4][i-i0] = reg_cons_1_p1;
    reg_cons_1_p1 = reg_cons_1_p2;
    reg_cons_1_p2 = reg_cons_1_p3;
    reg_cons_1_p3 = reg_cons_1_p4;
    reg_cons_2_m4 = reg_cons_2_m3;
    reg_cons_2_m3 = reg_cons_2_m2;
    reg_cons_2_m2 = reg_cons_2_m1;
    reg_cons_2_m1 = sh_cons_2_c0[j-j0+4][i-i0];
    sh_cons_2_c0[j-j0+4][i-i0] = reg_cons_2_p1;
    reg_cons_2_p1 = reg_cons_2_p2;
    reg_cons_2_p2 = reg_cons_2_p3;
    reg_cons_2_p3 = reg_cons_2_p4;
    reg_cons_3_m4 = reg_cons_3_m3;
    reg_cons_3_m3 = reg_cons_3_m2;
    reg_cons_3_m2 = reg_cons_3_m1;
    reg_cons_3_m1 = sh_cons_3_c0[j-j0+4][i-i0];
    sh_cons_3_c0[j-j0+4][i-i0] = reg_cons_3_p1;
    reg_cons_3_p1 = reg_cons_3_p2;
    reg_cons_3_p2 = reg_cons_3_p3;
    reg_cons_3_p3 = reg_cons_3_p4;
    reg_cons_4_m4 = reg_cons_4_m3;
    reg_cons_4_m3 = reg_cons_4_m2;
    reg_cons_4_m2 = reg_cons_4_m1;
    reg_cons_4_m1 = sh_cons_4_c0[j-j0+4][i-i0];
    sh_cons_4_c0[j-j0+4][i-i0] = reg_cons_4_p1;
    reg_cons_4_p1 = reg_cons_4_p2;
    reg_cons_4_p2 = reg_cons_4_p3;
    reg_cons_4_p3 = reg_cons_4_p4;
    reg_q_3_m4 = reg_q_3_m3;
    reg_q_3_m3 = reg_q_3_m2;
    reg_q_3_m2 = reg_q_3_m1;
    reg_q_3_m1 = reg_q_3_c0;
    reg_q_3_c0 = reg_q_3_p1;
    reg_q_3_p1 = reg_q_3_p2;
    reg_q_3_p2 = reg_q_3_p3;
    reg_q_3_p3 = reg_q_3_p4;
    reg_q_4_m4 = reg_q_4_m3;
    reg_q_4_m3 = reg_q_4_m2;
    reg_q_4_m2 = reg_q_4_m1;
    reg_q_4_m1 = sh_q_4_c0[j-j0+4][i-i0];
    sh_q_4_c0[j-j0+4][i-i0] = reg_q_4_p1;
    reg_q_4_p1 = reg_q_4_p2;
    reg_q_4_p2 = reg_q_4_p3;
    reg_q_4_p3 = reg_q_4_p4;
   }
  }
}

extern "C" void host_code (double *h_flux_0, double *h_flux_1, double *h_flux_2, double *h_flux_3, double *h_flux_4, double *h_cons_1, double *h_cons_2, double *h_cons_3, double *h_cons_4, double *h_q_1, double *h_q_2, double *h_q_3, double *h_q_4, double dxinv0, double dxinv1, double dxinv2, int L, int M, int N) {
  double *flux_0;
  cudaMalloc (&flux_0, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for flux_0\n");
  cudaMemcpy (flux_0, h_flux_0, sizeof(double)*L*M*N, cudaMemcpyHostToDevice);
  double *flux_1;
  cudaMalloc (&flux_1, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for flux_1\n");
  cudaMemcpy (flux_1, h_flux_1, sizeof(double)*L*M*N, cudaMemcpyHostToDevice);
  double *flux_2;
  cudaMalloc (&flux_2, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for flux_2\n");
  cudaMemcpy (flux_2, h_flux_2, sizeof(double)*L*M*N, cudaMemcpyHostToDevice);
  double *flux_3;
  cudaMalloc (&flux_3, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for flux_3\n");
  cudaMemcpy (flux_3, h_flux_3, sizeof(double)*L*M*N, cudaMemcpyHostToDevice);
  double *flux_4;
  cudaMalloc (&flux_4, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for flux_4\n");
  cudaMemcpy (flux_4, h_flux_4, sizeof(double)*L*M*N, cudaMemcpyHostToDevice);
  double *cons_1;
  cudaMalloc (&cons_1, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for cons_1\n");
  cudaMemcpy (cons_1, h_cons_1, sizeof(double)*L*M*N, cudaMemcpyHostToDevice);
  double *cons_2;
  cudaMalloc (&cons_2, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for cons_2\n");
  cudaMemcpy (cons_2, h_cons_2, sizeof(double)*L*M*N, cudaMemcpyHostToDevice);
  double *cons_3;
  cudaMalloc (&cons_3, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for cons_3\n");
  cudaMemcpy (cons_3, h_cons_3, sizeof(double)*L*M*N, cudaMemcpyHostToDevice);
  double *cons_4;
  cudaMalloc (&cons_4, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for cons_4\n");
  cudaMemcpy (cons_4, h_cons_4, sizeof(double)*L*M*N, cudaMemcpyHostToDevice);
  double *q_1;
  cudaMalloc (&q_1, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for q_1\n");
  cudaMemcpy (q_1, h_q_1, sizeof(double)*L*M*N, cudaMemcpyHostToDevice);
  double *q_2;
  cudaMalloc (&q_2, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for q_2\n");
  cudaMemcpy (q_2, h_q_2, sizeof(double)*L*M*N, cudaMemcpyHostToDevice);
  double *q_3;
  cudaMalloc (&q_3, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for q_3\n");
  cudaMemcpy (q_3, h_q_3, sizeof(double)*L*M*N, cudaMemcpyHostToDevice);
  double *q_4;
  cudaMalloc (&q_4, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for q_4\n");
  cudaMemcpy (q_4, h_q_4, sizeof(double)*L*M*N, cudaMemcpyHostToDevice);

  dim3 blockconfig_1 (16, 16, 1);
  dim3 gridconfig_1 (ceil(N, blockconfig_1.x-8), ceil(M-8, blockconfig_1.y), 1);
  hypterm <<<gridconfig_1, blockconfig_1>>> (flux_0, flux_1, flux_2, flux_3, flux_4, cons_1, cons_2, cons_3, cons_4, q_1, q_2, q_3, q_4, dxinv0, dxinv1, dxinv2, L, M, N);

  cudaMemcpy (h_flux_0, flux_0, sizeof(double)*L*M*N, cudaMemcpyDeviceToHost);
  cudaMemcpy (h_flux_1, flux_1, sizeof(double)*L*M*N, cudaMemcpyDeviceToHost);
  cudaMemcpy (h_flux_2, flux_2, sizeof(double)*L*M*N, cudaMemcpyDeviceToHost);
  cudaMemcpy (h_flux_3, flux_3, sizeof(double)*L*M*N, cudaMemcpyDeviceToHost);
  cudaMemcpy (h_flux_4, flux_4, sizeof(double)*L*M*N, cudaMemcpyDeviceToHost);

  //Free allocated memory
  cudaFree (flux_0);
  cudaFree (flux_1);
  cudaFree (flux_2);
  cudaFree (flux_3);
  cudaFree (flux_4);
  cudaFree (cons_1);
  cudaFree (cons_2);
  cudaFree (cons_3);
  cudaFree (cons_4);
  cudaFree (q_1);
  cudaFree (q_2);
  cudaFree (q_3);
  cudaFree (q_4);
}

	.file	"tmpxft_0039a1d2_00000000-6_hypterm-halo.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2031:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2031:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0, @function
_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0:
.LFB2173:
	.cfi_startproc
	jmp	cudaMalloc@PLT
	.cfi_endproc
.LFE2173:
	.size	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0, .-_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"CUDA error : %s, %s\n"
	.text
	.globl	_Z11check_errorPKc
	.type	_Z11check_errorPKc, @function
_Z11check_errorPKc:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rdi, %rbx
	call	cudaGetLastError@PLT
	testl	%eax, %eax
	je	.L3
	movl	%eax, %edi
	call	cudaGetErrorString@PLT
	movl	$2, %edi
	movq	%rbx, %rdx
	leaq	.LC0(%rip), %rsi
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__printf_chk@PLT
	orl	$-1, %edi
	call	exit@PLT
.L3:
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z11check_errorPKc, .-_Z11check_errorPKc
	.globl	_Z56__device_stub__Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiiiPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii
	.type	_Z56__device_stub__Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiiiPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii, @function
_Z56__device_stub__Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiiiPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii:
.LFB2053:
	.cfi_startproc
	endbr64
	subq	$360, %rsp
	.cfi_def_cfa_offset 368
	movq	368(%rsp), %rax
	movq	%rdi, 120(%rsp)
	leaq	144(%rsp), %rdi
	movq	%rsi, 112(%rsp)
	leaq	156(%rsp), %rsi
	movq	%rax, 72(%rsp)
	movq	376(%rsp), %rax
	movq	%rdx, 104(%rsp)
	leaq	128(%rsp), %rdx
	movq	%rax, 64(%rsp)
	movq	384(%rsp), %rax
	movq	%rcx, 96(%rsp)
	leaq	136(%rsp), %rcx
	movq	%rax, 56(%rsp)
	movq	392(%rsp), %rax
	movq	%r8, 88(%rsp)
	movq	%rax, 48(%rsp)
	movq	400(%rsp), %rax
	movq	%r9, 80(%rsp)
	movq	%rax, 40(%rsp)
	movq	408(%rsp), %rax
	movsd	%xmm0, 16(%rsp)
	movq	%rax, 32(%rsp)
	movq	416(%rsp), %rax
	movsd	%xmm1, 8(%rsp)
	movq	%rax, 24(%rsp)
	movsd	%xmm2, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 344(%rsp)
	xorl	%eax, %eax
	leaq	120(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	112(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	104(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	96(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	88(%rsp), %rax
	movq	%rax, 224(%rsp)
	leaq	80(%rsp), %rax
	movq	%rax, 232(%rsp)
	leaq	72(%rsp), %rax
	movq	%rax, 240(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 248(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 256(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 264(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 272(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 280(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 288(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 296(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 304(%rsp)
	movq	%rsp, %rax
	movq	%rax, 312(%rsp)
	leaq	424(%rsp), %rax
	movq	%rax, 320(%rsp)
	leaq	432(%rsp), %rax
	movq	%rax, 328(%rsp)
	leaq	440(%rsp), %rax
	movq	%rax, 336(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 144(%rsp)
	movl	$1, 152(%rsp)
	movq	%rax, 156(%rsp)
	movl	$1, 164(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L7
	pushq	136(%rsp)
	.cfi_def_cfa_offset 376
	leaq	_Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii(%rip), %rdi
	pushq	136(%rsp)
	.cfi_def_cfa_offset 384
	movq	172(%rsp), %rcx
	movl	180(%rsp), %r8d
	movq	160(%rsp), %rsi
	movl	168(%rsp), %edx
	leaq	208(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 376
	popq	%rdx
	.cfi_def_cfa_offset 368
.L7:
	movq	344(%rsp), %rax
	subq	%fs:40, %rax
	je	.L9
	call	__stack_chk_fail@PLT
.L9:
	addq	$360, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2053:
	.size	_Z56__device_stub__Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiiiPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii, .-_Z56__device_stub__Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiiiPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii
	.globl	_Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii
	.type	_Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii, @function
_Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii:
.LFB2054:
	.cfi_startproc
	endbr64
	jmp	_Z56__device_stub__Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiiiPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii
	.cfi_endproc
.LFE2054:
	.size	_Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii, .-_Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii
	.section	.rodata.str1.1
.LC1:
	.string	"Failed to allocate device memory for flux_0\n"
.LC2:
	.string	"Failed to allocate device memory for flux_1\n"
.LC3:
	.string	"Failed to allocate device memory for flux_2\n"
.LC4:
	.string	"Failed to allocate device memory for flux_3\n"
.LC5:
	.string	"Failed to allocate device memory for flux_4\n"
.LC6:
	.string	"Failed to allocate device memory for cons_1\n"
.LC7:
	.string	"Failed to allocate device memory for cons_2\n"
.LC8:
	.string	"Failed to allocate device memory for cons_3\n"
.LC9:
	.string	"Failed to allocate device memory for cons_4\n"
.LC10:
	.string	"Failed to allocate device memory for q_1\n"
.LC11:
	.string	"Failed to allocate device memory for q_2\n"
.LC12:
	.string	"Failed to allocate device memory for q_3\n"
.LC13:
	.string	"Failed to allocate device memory for q_4\n"
	.text
	.globl	host_code
	.type	host_code, @function
host_code:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%r9, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$264, %rsp
	.cfi_def_cfa_offset 320
	movq	344(%rsp), %rax
	movq	%r8, 48(%rsp)
	movsd	%xmm0, 104(%rsp)
	movq	320(%rsp), %r15
	movq	328(%rsp), %r14
	movq	%rax, 56(%rsp)
	movq	352(%rsp), %rax
	movsd	%xmm1, 96(%rsp)
	movq	336(%rsp), %r13
	movq	%rax, 64(%rsp)
	movq	360(%rsp), %rax
	movq	%rcx, 40(%rsp)
	movq	%rax, 72(%rsp)
	movq	368(%rsp), %rax
	movq	%rdi, 16(%rsp)
	leaq	120(%rsp), %rdi
	movq	%rsi, 24(%rsp)
	movq	%rdx, 32(%rsp)
	movq	%rax, 80(%rsp)
	movsd	%xmm2, 88(%rsp)
	movl	384(%rsp), %ebp
	movslq	392(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 248(%rsp)
	xorl	%eax, %eax
	movslq	%ebp, %rax
	movq	%rdx, %rbx
	imulq	%rdx, %rax
	movslq	376(%rsp), %rdx
	imulq	%rdx, %rax
	salq	$3, %rax
	movq	%rax, %rsi
	movq	%rax, 8(%rsp)
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	leaq	.LC1(%rip), %rdi
	call	_Z11check_errorPKc
	movq	8(%rsp), %rdx
	movq	16(%rsp), %rsi
	movl	$1, %ecx
	movq	120(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rsi
	leaq	128(%rsp), %rdi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	leaq	.LC2(%rip), %rdi
	call	_Z11check_errorPKc
	movq	8(%rsp), %rdx
	movq	24(%rsp), %rsi
	movl	$1, %ecx
	movq	128(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rsi
	leaq	136(%rsp), %rdi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	leaq	.LC3(%rip), %rdi
	call	_Z11check_errorPKc
	movq	8(%rsp), %rdx
	movq	32(%rsp), %rsi
	movl	$1, %ecx
	movq	136(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rsi
	leaq	144(%rsp), %rdi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	leaq	.LC4(%rip), %rdi
	call	_Z11check_errorPKc
	movq	8(%rsp), %rdx
	movq	40(%rsp), %rsi
	movl	$1, %ecx
	movq	144(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rsi
	leaq	152(%rsp), %rdi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	leaq	.LC5(%rip), %rdi
	call	_Z11check_errorPKc
	movq	8(%rsp), %rdx
	movq	48(%rsp), %rsi
	movl	$1, %ecx
	movq	152(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rsi
	leaq	160(%rsp), %rdi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	leaq	.LC6(%rip), %rdi
	call	_Z11check_errorPKc
	movq	8(%rsp), %rdx
	movl	$1, %ecx
	movq	%r12, %rsi
	movq	160(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rsi
	leaq	168(%rsp), %rdi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	leaq	.LC7(%rip), %rdi
	call	_Z11check_errorPKc
	movq	8(%rsp), %rdx
	movl	$1, %ecx
	movq	%r15, %rsi
	movq	168(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rsi
	leaq	176(%rsp), %rdi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	leaq	.LC8(%rip), %rdi
	call	_Z11check_errorPKc
	movq	8(%rsp), %rdx
	movl	$1, %ecx
	movq	%r14, %rsi
	movq	176(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rsi
	leaq	184(%rsp), %rdi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	leaq	.LC9(%rip), %rdi
	call	_Z11check_errorPKc
	movq	8(%rsp), %rdx
	movl	$1, %ecx
	movq	%r13, %rsi
	movq	184(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rsi
	leaq	192(%rsp), %rdi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	leaq	.LC10(%rip), %rdi
	call	_Z11check_errorPKc
	movq	8(%rsp), %rdx
	movq	56(%rsp), %rsi
	movl	$1, %ecx
	movq	192(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rsi
	leaq	200(%rsp), %rdi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	leaq	.LC11(%rip), %rdi
	call	_Z11check_errorPKc
	movq	8(%rsp), %rdx
	movq	64(%rsp), %rsi
	movl	$1, %ecx
	movq	200(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rsi
	leaq	208(%rsp), %rdi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	leaq	.LC12(%rip), %rdi
	call	_Z11check_errorPKc
	movq	8(%rsp), %rdx
	movq	72(%rsp), %rsi
	movl	$1, %ecx
	movq	208(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rsi
	leaq	216(%rsp), %rdi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	leaq	.LC13(%rip), %rdi
	call	_Z11check_errorPKc
	movq	8(%rsp), %rdx
	movq	80(%rsp), %rsi
	movl	$1, %ecx
	movq	216(%rsp), %rdi
	call	cudaMemcpy@PLT
	leal	-8(%rbp), %eax
	movl	%ebx, %ecx
	movl	$1, %esi
	movl	%eax, %edx
	andl	$15, %eax
	shrl	$4, %edx
	cmpl	$1, %eax
	movl	%ebx, %eax
	sbbl	$-1, %edx
	andl	$7, %ecx
	shrl	$3, %eax
	cmpl	$1, %ecx
	movl	%edx, 240(%rsp)
	movl	$1, %ecx
	movabsq	$68719476752, %rdx
	sbbl	$-1, %eax
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	%eax, 236(%rsp)
	movq	236(%rsp), %rdi
	call	__cudaPushCallConfiguration@PLT
	movsd	88(%rsp), %xmm2
	movsd	96(%rsp), %xmm1
	testl	%eax, %eax
	movsd	104(%rsp), %xmm0
	jne	.L15
	pushq	%rbx
	.cfi_def_cfa_offset 328
	pushq	%rbp
	.cfi_def_cfa_offset 336
	movl	392(%rsp), %eax
	pushq	%rax
	.cfi_def_cfa_offset 344
	pushq	240(%rsp)
	.cfi_def_cfa_offset 352
	pushq	240(%rsp)
	.cfi_def_cfa_offset 360
	pushq	240(%rsp)
	.cfi_def_cfa_offset 368
	pushq	240(%rsp)
	.cfi_def_cfa_offset 376
	pushq	240(%rsp)
	.cfi_def_cfa_offset 384
	pushq	240(%rsp)
	.cfi_def_cfa_offset 392
	pushq	240(%rsp)
	.cfi_def_cfa_offset 400
	movq	240(%rsp), %r9
	movq	232(%rsp), %r8
	movq	224(%rsp), %rcx
	movq	216(%rsp), %rdx
	movq	208(%rsp), %rsi
	movq	200(%rsp), %rdi
	call	_Z56__device_stub__Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiiiPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii
	addq	$80, %rsp
	.cfi_def_cfa_offset 320
.L15:
	movq	8(%rsp), %rdx
	movq	120(%rsp), %rsi
	movl	$2, %ecx
	movq	16(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rdx
	movq	24(%rsp), %rdi
	movl	$2, %ecx
	movq	128(%rsp), %rsi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rdx
	movq	32(%rsp), %rdi
	movl	$2, %ecx
	movq	136(%rsp), %rsi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rdx
	movq	40(%rsp), %rdi
	movl	$2, %ecx
	movq	144(%rsp), %rsi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rdx
	movq	48(%rsp), %rdi
	movl	$2, %ecx
	movq	152(%rsp), %rsi
	call	cudaMemcpy@PLT
	movq	120(%rsp), %rdi
	call	cudaFree@PLT
	movq	128(%rsp), %rdi
	call	cudaFree@PLT
	movq	136(%rsp), %rdi
	call	cudaFree@PLT
	movq	144(%rsp), %rdi
	call	cudaFree@PLT
	movq	152(%rsp), %rdi
	call	cudaFree@PLT
	movq	160(%rsp), %rdi
	call	cudaFree@PLT
	movq	168(%rsp), %rdi
	call	cudaFree@PLT
	movq	176(%rsp), %rdi
	call	cudaFree@PLT
	movq	184(%rsp), %rdi
	call	cudaFree@PLT
	movq	192(%rsp), %rdi
	call	cudaFree@PLT
	movq	200(%rsp), %rdi
	call	cudaFree@PLT
	movq	208(%rsp), %rdi
	call	cudaFree@PLT
	movq	216(%rsp), %rdi
	call	cudaFree@PLT
	movq	248(%rsp), %rax
	subq	%fs:40, %rax
	je	.L16
	call	__stack_chk_fail@PLT
.L16:
	addq	$264, %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
.LFE2028:
	.size	host_code, .-host_code
	.section	.rodata.str1.1
.LC14:
	.string	"_Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2056:
	.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	.LC14(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii(%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
.LFE2056:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

#include <stdio.h>
#include "hip/hip_runtime.h"
#define max(x,y)  ((x) > (y)? (x) : (y))
#define min(x,y)  ((x) < (y)? (x) : (y))
#define ceil(a,b) ((a) % (b) == 0 ? (a) / (b) : ((a) / (b)) + 1)

void check_error (const char* message) {
  hipError_t error = hipGetLastError ();
  if (error != hipSuccess) {
    printf ("CUDA error : %s, %s\n", message, hipGetErrorString (error));
    exit(-1);
  }
}

__global__ void hypterm (double * flux_0, double * flux_1, double * flux_2, double * flux_3, double * flux_4, double * cons_1, double * cons_2, double * cons_3, double * cons_4, double * q_1, double * q_2, double * q_3, double * q_4, double dxinv0, double dxinv1, double dxinv2, int L, int M, int N) {
  //Determining the block's indices
  int blockdim_i= (int)(blockDim.x);
  int i0 = (int)(blockIdx.x)*(blockdim_i-8);
  int i = i0 + (int)(threadIdx.x);
  int blockdim_j= (int)(blockDim.y);
  int j0 = (int)(blockIdx.y)*(blockdim_j)+4;
  int j = j0 + (int)(threadIdx.y);

  //Declarations
  double reg_cons_1_m4=0, reg_cons_1_m3=0, reg_cons_1_m2=0, reg_cons_1_m1=0, __shared__ sh_cons_1_c0[24][16], reg_cons_1_p1=0, reg_cons_1_p2=0, reg_cons_1_p3=0, reg_cons_1_p4=0;
  double reg_cons_2_m4=0, reg_cons_2_m3=0, reg_cons_2_m2=0, reg_cons_2_m1=0, __shared__ sh_cons_2_c0[24][16], reg_cons_2_p1=0, reg_cons_2_p2=0, reg_cons_2_p3=0, reg_cons_2_p4=0;
  double reg_cons_3_m4=0, reg_cons_3_m3=0, reg_cons_3_m2=0, reg_cons_3_m1=0, __shared__ sh_cons_3_c0[24][16], reg_cons_3_p1=0, reg_cons_3_p2=0, reg_cons_3_p3=0, reg_cons_3_p4=0;
  double reg_cons_4_m4=0, reg_cons_4_m3=0, reg_cons_4_m2=0, reg_cons_4_m1=0, __shared__ sh_cons_4_c0[24][16], reg_cons_4_p1=0, reg_cons_4_p2=0, reg_cons_4_p3=0, reg_cons_4_p4=0;
  double __shared__ sh_q_1_c0[24][16];
  double __shared__ sh_q_2_c0[24][16];
  double reg_q_3_m4=0, reg_q_3_m3=0, reg_q_3_m2=0, reg_q_3_m1=0, reg_q_3_c0=0, reg_q_3_p1=0, reg_q_3_p2=0, reg_q_3_p3=0, reg_q_3_p4=0;
  double reg_q_4_m4=0, reg_q_4_m3=0, reg_q_4_m2=0, reg_q_4_m1=0, __shared__ sh_q_4_c0[24][16], reg_q_4_p1=0, reg_q_4_p2=0, reg_q_4_p3=0, reg_q_4_p4=0;

  //Value Initialization
  if (j <= M-1 & i <= min (i0+blockdim_i-1, N-1)) {
    reg_cons_1_m4 = cons_1[0 + j*N + i];
    reg_cons_1_m3 = cons_1[1*M*N + j*N + i];
    reg_cons_1_m2 = cons_1[2*M*N + j*N + i];
    reg_cons_1_m1 = cons_1[3*M*N + j*N + i];
    sh_cons_1_c0[j-j0+4][i-i0] = cons_1[4*M*N + j*N + i];
    reg_cons_1_p1 = cons_1[5*M*N + j*N + i];
    reg_cons_1_p2 = cons_1[6*M*N + j*N + i];
    reg_cons_1_p3 = cons_1[7*M*N + j*N + i];
    reg_cons_2_m4 = cons_2[0 + j*N + i];
    reg_cons_2_m3 = cons_2[1*M*N + j*N + i];
    reg_cons_2_m2 = cons_2[2*M*N + j*N + i];
    reg_cons_2_m1 = cons_2[3*M*N + j*N + i];
    sh_cons_2_c0[j-j0+4][i-i0] = cons_2[4*M*N + j*N + i];
    reg_cons_2_p1 = cons_2[5*M*N + j*N + i];
    reg_cons_2_p2 = cons_2[6*M*N + j*N + i];
    reg_cons_2_p3 = cons_2[7*M*N + j*N + i];
    reg_cons_3_m4 = cons_3[0 + j*N + i];
    reg_cons_3_m3 = cons_3[1*M*N + j*N + i];
    reg_cons_3_m2 = cons_3[2*M*N + j*N + i];
    reg_cons_3_m1 = cons_3[3*M*N + j*N + i];
    sh_cons_3_c0[j-j0+4][i-i0] = cons_3[4*M*N + j*N + i];
    reg_cons_3_p1 = cons_3[5*M*N + j*N + i];
    reg_cons_3_p2 = cons_3[6*M*N + j*N + i];
    reg_cons_3_p3 = cons_3[7*M*N + j*N + i];
    reg_cons_4_m4 = cons_4[0 + j*N + i];
    reg_cons_4_m3 = cons_4[1*M*N + j*N + i];
    reg_cons_4_m2 = cons_4[2*M*N + j*N + i];
    reg_cons_4_m1 = cons_4[3*M*N + j*N + i];
    sh_cons_4_c0[j-j0+4][i-i0] = cons_4[4*M*N + j*N + i];
    reg_cons_4_p1 = cons_4[5*M*N + j*N + i];
    reg_cons_4_p2 = cons_4[6*M*N + j*N + i];
    reg_cons_4_p3 = cons_4[7*M*N + j*N + i];
    reg_q_3_m4 = q_3[0 + j*N + i];
    reg_q_3_m3 = q_3[1*M*N + j*N + i];
    reg_q_3_m2 = q_3[2*M*N + j*N + i];
    reg_q_3_m1 = q_3[3*M*N + j*N + i];
    reg_q_3_c0 = q_3[4*M*N + j*N + i];
    reg_q_3_p1 = q_3[5*M*N + j*N + i];
    reg_q_3_p2 = q_3[6*M*N + j*N + i];
    reg_q_3_p3 = q_3[7*M*N + j*N + i];
    reg_q_4_m4 = q_4[0 + j*N + i];
    reg_q_4_m3 = q_4[1*M*N + j*N + i];
    reg_q_4_m2 = q_4[2*M*N + j*N + i];
    reg_q_4_m1 = q_4[3*M*N + j*N + i];
    sh_q_4_c0[j-j0+4][i-i0] = q_4[4*M*N + j*N + i];
    reg_q_4_p1 = q_4[5*M*N + j*N + i];
    reg_q_4_p2 = q_4[6*M*N + j*N + i];
    reg_q_4_p3 = q_4[7*M*N + j*N + i];
  }
  if(threadIdx.y < 4) {
   int jj = (j-4);
   if (jj <= M-1 & i <= min (i0+blockdim_i-1, N-1)) {
    sh_cons_1_c0[jj-j0+4][i-i0] = cons_1[4*M*N + jj*N + i];
    sh_cons_2_c0[jj-j0+4][i-i0] = cons_2[4*M*N + jj*N + i];
    sh_cons_3_c0[jj-j0+4][i-i0] = cons_3[4*M*N + jj*N + i];
    sh_cons_4_c0[jj-j0+4][i-i0] = cons_4[4*M*N + jj*N + i];
    sh_q_4_c0[jj-j0+4][i-i0] = q_4[4*M*N + jj*N + i];
   }
  } else if(threadIdx.y < 8) {
   int jj = (j-4)+16;
   if (jj <= M-1 & i <= min (i0+blockdim_i-1, N-1)) {
    sh_cons_1_c0[jj-j0+4][i-i0] = cons_1[4*M*N + jj*N + i];
    sh_cons_2_c0[jj-j0+4][i-i0] = cons_2[4*M*N + jj*N + i];
    sh_cons_3_c0[jj-j0+4][i-i0] = cons_3[4*M*N + jj*N + i];
    sh_cons_4_c0[jj-j0+4][i-i0] = cons_4[4*M*N + jj*N + i];
    sh_q_4_c0[jj-j0+4][i-i0] = q_4[4*M*N + jj*N + i];
   }
  }

  //Rest of the computation
  for (int k=4; k<=L-5; ++k) {
    //Fetch new plane
    if (j <= M-1 & i <= min (i0+blockdim_i-1, N-1)) {
        reg_cons_1_p4 = cons_1[(k+4)*M*N + j*N + i];
        reg_cons_2_p4 = cons_2[(k+4)*M*N + j*N + i];
        reg_cons_3_p4 = cons_3[(k+4)*M*N + j*N + i];
        reg_cons_4_p4 = cons_4[(k+4)*M*N + j*N + i];
        sh_q_1_c0[j-j0+4][i-i0] = q_1[k*M*N + j*N + i];
        sh_q_2_c0[j-j0+4][i-i0] = q_2[k*M*N + j*N + i];
        reg_q_3_p4 = q_3[(k+4)*M*N + j*N + i];
        reg_q_4_p4 = q_4[(k+4)*M*N + j*N + i];
    }
    if(threadIdx.y < 4) {
     int jj = (j-4);
     if (jj <= M-1 & i <= min (i0+blockdim_i-1, N-1)) {
        sh_cons_1_c0[jj-j0+4][i-i0] = cons_1[k*M*N + jj*N + i];
        sh_cons_2_c0[jj-j0+4][i-i0] = cons_2[k*M*N + jj*N + i];
        sh_cons_3_c0[jj-j0+4][i-i0] = cons_3[k*M*N + jj*N + i];
        sh_cons_4_c0[jj-j0+4][i-i0] = cons_4[k*M*N + jj*N + i];
     }
    } else if(threadIdx.y < 8) {
     int jj = (j-4)+16;
     if (jj <= M-1 & i <= min (i0+blockdim_i-1, N-1)) {
        sh_cons_1_c0[jj-j0+4][i-i0] = cons_1[k*M*N + jj*N + i];
        sh_cons_2_c0[jj-j0+4][i-i0] = cons_2[k*M*N + jj*N + i];
        sh_cons_3_c0[jj-j0+4][i-i0] = cons_3[k*M*N + jj*N + i];
        sh_cons_4_c0[jj-j0+4][i-i0] = cons_4[k*M*N + jj*N + i];
     }
    } else if(threadIdx.y < 12) {
     int jj = (j-12);
     if (jj <= M-1 & i <= min (i0+blockdim_i-1, N-1)) {
        sh_q_1_c0[jj-j0+4][i-i0] = q_1[k*M*N + jj*N + i];
        sh_q_2_c0[jj-j0+4][i-i0] = q_2[k*M*N + jj*N + i];
        sh_q_4_c0[jj-j0+4][i-i0] = q_4[k*M*N + jj*N + i];
     }
    } else {
     int jj = (j-12)+16;
     if (jj <= M-1 & i <= min (i0+blockdim_i-1, N-1)) {
        sh_q_1_c0[jj-j0+4][i-i0] = q_1[k*M*N + jj*N + i];
        sh_q_2_c0[jj-j0+4][i-i0] = q_2[k*M*N + jj*N + i];
        sh_q_4_c0[jj-j0+4][i-i0] = q_4[k*M*N + jj*N + i];
     }
    }
    __syncthreads ();
	double r0,r1,r2,r3,r4;
    if ((j <= M-5) & i >= max (i0+4, 4) & i <= min (i0+blockdim_i-5, N-5)) {
	r0=flux_0[k*M*N + j*N + i], r1=flux_1[k*M*N + j*N + i], r2=flux_2[k*M*N + j*N + i], r3=flux_3[k*M*N + j*N + i], r4 = flux_4[k*M*N + j*N + i];
//	double r0=0.0f,r1=0.0f,r2=0.0f,r3=0.0f,r4=0.0f;
        r0 -= (((((0.8f * (sh_cons_1_c0[j-j0+4][i-i0+1] - sh_cons_1_c0[j-j0+4][i-i0-1])) - (0.2f * (sh_cons_1_c0[j-j0+4][i-i0+2] - sh_cons_1_c0[j-j0+4][i-i0-2]))) + (0.038f * (sh_cons_1_c0[j-j0+4][i-i0+3] - sh_cons_1_c0[j-j0+4][i-i0-3]))) - (0.0035f * (sh_cons_1_c0[j-j0+4][i-i0+4] - sh_cons_1_c0[j-j0+4][i-i0-4]))) * dxinv0);
        r0 -= (((((0.8f * (sh_cons_2_c0[j-j0+4+1][i-i0] - sh_cons_2_c0[j-j0+4-1][i-i0])) - (0.2f * (sh_cons_2_c0[j-j0+4+2][i-i0] - sh_cons_2_c0[j-j0+4-2][i-i0]))) + (0.038f * (sh_cons_2_c0[j-j0+4+3][i-i0] - sh_cons_2_c0[j-j0+4-3][i-i0]))) - (0.0035f * (sh_cons_2_c0[j-j0+4+4][i-i0] - sh_cons_2_c0[j-j0+4-4][i-i0]))) * dxinv1);
        r0 -= (((((0.8f * (reg_cons_3_p1 - reg_cons_3_m1)) - (0.2f * (reg_cons_3_p2 - reg_cons_3_m2))) + (0.038f * (reg_cons_3_p3 - reg_cons_3_m3))) - (0.0035f * (reg_cons_3_p4 - reg_cons_3_m4))) * dxinv2);
	flux_0[k*M*N + j*N + i] = r0;
        r1 -= (((((0.8f * (((sh_cons_1_c0[j-j0+4][i-i0+1] * sh_q_1_c0[j-j0+4][i-i0+1]) - (sh_cons_1_c0[j-j0+4][i-i0-1] * sh_q_1_c0[j-j0+4][i-i0-1])) + (sh_q_4_c0[j-j0+4][i-i0+1] - sh_q_4_c0[j-j0+4][i-i0-1]))) - (0.2f * (((sh_cons_1_c0[j-j0+4][i-i0+2] * sh_q_1_c0[j-j0+4][i-i0+2]) - (sh_cons_1_c0[j-j0+4][i-i0-2] * sh_q_1_c0[j-j0+4][i-i0-2])) + (sh_q_4_c0[j-j0+4][i-i0+2] - sh_q_4_c0[j-j0+4][i-i0-2])))) + (0.038f * (((sh_cons_1_c0[j-j0+4][i-i0+3] * sh_q_1_c0[j-j0+4][i-i0+3]) - (sh_cons_1_c0[j-j0+4][i-i0-3] * sh_q_1_c0[j-j0+4][i-i0-3])) + (sh_q_4_c0[j-j0+4][i-i0+3] - sh_q_4_c0[j-j0+4][i-i0-3])))) - (0.0035f * (((sh_cons_1_c0[j-j0+4][i-i0+4] * sh_q_1_c0[j-j0+4][i-i0+4]) - (sh_cons_1_c0[j-j0+4][i-i0-4] * sh_q_1_c0[j-j0+4][i-i0-4])) + (sh_q_4_c0[j-j0+4][i-i0+4] - sh_q_4_c0[j-j0+4][i-i0-4])))) * dxinv0);
        r1 -= (((((0.8f * ((sh_cons_1_c0[j-j0+4+1][i-i0] * sh_q_2_c0[j-j0+4+1][i-i0]) - (sh_cons_1_c0[j-j0+4-1][i-i0] * sh_q_2_c0[j-j0+4-1][i-i0]))) - (0.2f * ((sh_cons_1_c0[j-j0+4+2][i-i0] * sh_q_2_c0[j-j0+4+2][i-i0]) - (sh_cons_1_c0[j-j0+4-2][i-i0] * sh_q_2_c0[j-j0+4-2][i-i0])))) + (0.038f * ((sh_cons_1_c0[j-j0+4+3][i-i0] * sh_q_2_c0[j-j0+4+3][i-i0]) - (sh_cons_1_c0[j-j0+4-3][i-i0] * sh_q_2_c0[j-j0+4-3][i-i0])))) - (0.0035f * ((sh_cons_1_c0[j-j0+4+4][i-i0] * sh_q_2_c0[j-j0+4+4][i-i0]) - (sh_cons_1_c0[j-j0+4-4][i-i0] * sh_q_2_c0[j-j0+4-4][i-i0])))) * dxinv1);
        r1 -= (((((0.8f * ((reg_cons_1_p1 * reg_q_3_p1) - (reg_cons_1_m1 * reg_q_3_m1))) - (0.2f * ((reg_cons_1_p2 * reg_q_3_p2) - (reg_cons_1_m2 * reg_q_3_m2)))) + (0.038f * ((reg_cons_1_p3 * reg_q_3_p3) - (reg_cons_1_m3 * reg_q_3_m3)))) - (0.0035f * ((reg_cons_1_p4 * reg_q_3_p4) - (reg_cons_1_m4 * reg_q_3_m4)))) * dxinv2);
	flux_1[k*M*N + j*N + i] = r1;
        r2 -= (((((0.8f * ((sh_cons_2_c0[j-j0+4][i-i0+1] * sh_q_1_c0[j-j0+4][i-i0+1]) - (sh_cons_2_c0[j-j0+4][i-i0-1] * sh_q_1_c0[j-j0+4][i-i0-1]))) - (0.2f * ((sh_cons_2_c0[j-j0+4][i-i0+2] * sh_q_1_c0[j-j0+4][i-i0+2]) - (sh_cons_2_c0[j-j0+4][i-i0-2] * sh_q_1_c0[j-j0+4][i-i0-2])))) + (0.038f * ((sh_cons_2_c0[j-j0+4][i-i0+3] * sh_q_1_c0[j-j0+4][i-i0+3]) - (sh_cons_2_c0[j-j0+4][i-i0-3] * sh_q_1_c0[j-j0+4][i-i0-3])))) - (0.0035f * ((sh_cons_2_c0[j-j0+4][i-i0+4] * sh_q_1_c0[j-j0+4][i-i0+4]) - (sh_cons_2_c0[j-j0+4][i-i0-4] * sh_q_1_c0[j-j0+4][i-i0-4])))) * dxinv0);
        r2 -= (((((0.8f * (((sh_cons_2_c0[j-j0+4+1][i-i0] * sh_q_2_c0[j-j0+4+1][i-i0]) - (sh_cons_2_c0[j-j0+4-1][i-i0] * sh_q_2_c0[j-j0+4-1][i-i0])) + (sh_q_4_c0[j-j0+4+1][i-i0] - sh_q_4_c0[j-j0+4-1][i-i0]))) - (0.2f * (((sh_cons_2_c0[j-j0+4+2][i-i0] * sh_q_2_c0[j-j0+4+2][i-i0]) - (sh_cons_2_c0[j-j0+4-2][i-i0] * sh_q_2_c0[j-j0+4-2][i-i0])) + (sh_q_4_c0[j-j0+4+2][i-i0] - sh_q_4_c0[j-j0+4-2][i-i0])))) + (0.038f * (((sh_cons_2_c0[j-j0+4+3][i-i0] * sh_q_2_c0[j-j0+4+3][i-i0]) - (sh_cons_2_c0[j-j0+4-3][i-i0] * sh_q_2_c0[j-j0+4-3][i-i0])) + (sh_q_4_c0[j-j0+4+3][i-i0] - sh_q_4_c0[j-j0+4-3][i-i0])))) - (0.0035f * (((sh_cons_2_c0[j-j0+4+4][i-i0] * sh_q_2_c0[j-j0+4+4][i-i0]) - (sh_cons_2_c0[j-j0+4-4][i-i0] * sh_q_2_c0[j-j0+4-4][i-i0])) + (sh_q_4_c0[j-j0+4+4][i-i0] - sh_q_4_c0[j-j0+4-4][i-i0])))) * dxinv1);
        r2 -= (((((0.8f * ((reg_cons_2_p1 * reg_q_3_p1) - (reg_cons_2_m1 * reg_q_3_m1))) - (0.2f * ((reg_cons_2_p2 * reg_q_3_p2) - (reg_cons_2_m2 * reg_q_3_m2)))) + (0.038f * ((reg_cons_2_p3 * reg_q_3_p3) - (reg_cons_2_m3 * reg_q_3_m3)))) - (0.0035f * ((reg_cons_2_p4 * reg_q_3_p4) - (reg_cons_2_m4 * reg_q_3_m4)))) * dxinv2);
	flux_2[k*M*N + j*N + i] = r2;
        r3 -= (((((0.8f * ((sh_cons_3_c0[j-j0+4][i-i0+1] * sh_q_1_c0[j-j0+4][i-i0+1]) - (sh_cons_3_c0[j-j0+4][i-i0-1] * sh_q_1_c0[j-j0+4][i-i0-1]))) - (0.2f * ((sh_cons_3_c0[j-j0+4][i-i0+2] * sh_q_1_c0[j-j0+4][i-i0+2]) - (sh_cons_3_c0[j-j0+4][i-i0-2] * sh_q_1_c0[j-j0+4][i-i0-2])))) + (0.038f * ((sh_cons_3_c0[j-j0+4][i-i0+3] * sh_q_1_c0[j-j0+4][i-i0+3]) - (sh_cons_3_c0[j-j0+4][i-i0-3] * sh_q_1_c0[j-j0+4][i-i0-3])))) - (0.0035f * ((sh_cons_3_c0[j-j0+4][i-i0+4] * sh_q_1_c0[j-j0+4][i-i0+4]) - (sh_cons_3_c0[j-j0+4][i-i0-4] * sh_q_1_c0[j-j0+4][i-i0-4])))) * dxinv0);
        r3 -= (((((0.8f * ((sh_cons_3_c0[j-j0+4+1][i-i0] * sh_q_2_c0[j-j0+4+1][i-i0]) - (sh_cons_3_c0[j-j0+4-1][i-i0] * sh_q_2_c0[j-j0+4-1][i-i0]))) - (0.2f * ((sh_cons_3_c0[j-j0+4+2][i-i0] * sh_q_2_c0[j-j0+4+2][i-i0]) - (sh_cons_3_c0[j-j0+4-2][i-i0] * sh_q_2_c0[j-j0+4-2][i-i0])))) + (0.038f * ((sh_cons_3_c0[j-j0+4+3][i-i0] * sh_q_2_c0[j-j0+4+3][i-i0]) - (sh_cons_3_c0[j-j0+4-3][i-i0] * sh_q_2_c0[j-j0+4-3][i-i0])))) - (0.0035f * ((sh_cons_3_c0[j-j0+4+4][i-i0] * sh_q_2_c0[j-j0+4+4][i-i0]) - (sh_cons_3_c0[j-j0+4-4][i-i0] * sh_q_2_c0[j-j0+4-4][i-i0])))) * dxinv1);
        r3 -= (((((0.8f * (((reg_cons_3_p1 * reg_q_3_p1) - (reg_cons_3_m1 * reg_q_3_m1)) + (reg_q_4_p1 - reg_q_4_m1))) - (0.2f * (((reg_cons_3_p2 * reg_q_3_p2) - (reg_cons_3_m2 * reg_q_3_m2)) + (reg_q_4_p2 - reg_q_4_m2)))) + (0.038f * (((reg_cons_3_p3 * reg_q_3_p3) - (reg_cons_3_m3 * reg_q_3_m3)) + (reg_q_4_p3 - reg_q_4_m3)))) - (0.0035f * (((reg_cons_3_p4 * reg_q_3_p4) - (reg_cons_3_m4 * reg_q_3_m4)) + (reg_q_4_p4 - reg_q_4_m4)))) * dxinv2);
	flux_3[k*M*N + j*N + i] = r3;
        r4 -= (((((0.8f * (((sh_cons_4_c0[j-j0+4][i-i0+1] * sh_q_1_c0[j-j0+4][i-i0+1]) - (sh_cons_4_c0[j-j0+4][i-i0-1] * sh_q_1_c0[j-j0+4][i-i0-1])) + ((sh_q_4_c0[j-j0+4][i-i0+1] * sh_q_1_c0[j-j0+4][i-i0+1]) - (sh_q_4_c0[j-j0+4][i-i0-1] * sh_q_1_c0[j-j0+4][i-i0-1])))) - (0.2f * (((sh_cons_4_c0[j-j0+4][i-i0+2] * sh_q_1_c0[j-j0+4][i-i0+2]) - (sh_cons_4_c0[j-j0+4][i-i0-2] * sh_q_1_c0[j-j0+4][i-i0-2])) + ((sh_q_4_c0[j-j0+4][i-i0+2] * sh_q_1_c0[j-j0+4][i-i0+2]) - (sh_q_4_c0[j-j0+4][i-i0-2] * sh_q_1_c0[j-j0+4][i-i0-2]))))) + (0.038f * (((sh_cons_4_c0[j-j0+4][i-i0+3] * sh_q_1_c0[j-j0+4][i-i0+3]) - (sh_cons_4_c0[j-j0+4][i-i0-3] * sh_q_1_c0[j-j0+4][i-i0-3])) + ((sh_q_4_c0[j-j0+4][i-i0+3] * sh_q_1_c0[j-j0+4][i-i0+3]) - (sh_q_4_c0[j-j0+4][i-i0-3] * sh_q_1_c0[j-j0+4][i-i0-3]))))) - (0.0035f * (((sh_cons_4_c0[j-j0+4][i-i0+4] * sh_q_1_c0[j-j0+4][i-i0+4]) - (sh_cons_4_c0[j-j0+4][i-i0-4] * sh_q_1_c0[j-j0+4][i-i0-4])) + ((sh_q_4_c0[j-j0+4][i-i0+4] * sh_q_1_c0[j-j0+4][i-i0+4]) - (sh_q_4_c0[j-j0+4][i-i0-4] * sh_q_1_c0[j-j0+4][i-i0-4]))))) * dxinv0);
        r4 -= (((((0.8f * (((sh_cons_4_c0[j-j0+4+1][i-i0] * sh_q_2_c0[j-j0+4+1][i-i0]) - (sh_cons_4_c0[j-j0+4-1][i-i0] * sh_q_2_c0[j-j0+4-1][i-i0])) + ((sh_q_4_c0[j-j0+4+1][i-i0] * sh_q_2_c0[j-j0+4+1][i-i0]) - (sh_q_4_c0[j-j0+4-1][i-i0] * sh_q_2_c0[j-j0+4-1][i-i0])))) - (0.2f * (((sh_cons_4_c0[j-j0+4+2][i-i0] * sh_q_2_c0[j-j0+4+2][i-i0]) - (sh_cons_4_c0[j-j0+4-2][i-i0] * sh_q_2_c0[j-j0+4-2][i-i0])) + ((sh_q_4_c0[j-j0+4+2][i-i0] * sh_q_2_c0[j-j0+4+2][i-i0]) - (sh_q_4_c0[j-j0+4-2][i-i0] * sh_q_2_c0[j-j0+4-2][i-i0]))))) + (0.038f * (((sh_cons_4_c0[j-j0+4+3][i-i0] * sh_q_2_c0[j-j0+4+3][i-i0]) - (sh_cons_4_c0[j-j0+4-3][i-i0] * sh_q_2_c0[j-j0+4-3][i-i0])) + ((sh_q_4_c0[j-j0+4+3][i-i0] * sh_q_2_c0[j-j0+4+3][i-i0]) - (sh_q_4_c0[j-j0+4-3][i-i0] * sh_q_2_c0[j-j0+4-3][i-i0]))))) - (0.0035f * (((sh_cons_4_c0[j-j0+4+4][i-i0] * sh_q_2_c0[j-j0+4+4][i-i0]) - (sh_cons_4_c0[j-j0+4-4][i-i0] * sh_q_2_c0[j-j0+4-4][i-i0])) + ((sh_q_4_c0[j-j0+4+4][i-i0] * sh_q_2_c0[j-j0+4+4][i-i0]) - (sh_q_4_c0[j-j0+4-4][i-i0] * sh_q_2_c0[j-j0+4-4][i-i0]))))) * dxinv1);
        r4 -= (((((0.8f * (((reg_cons_4_p1 * reg_q_3_p1) - (reg_cons_4_m1 * reg_q_3_m1)) + ((reg_q_4_p1 * reg_q_3_p1) - (reg_q_4_m1 * reg_q_3_m1)))) - (0.2f * (((reg_cons_4_p2 * reg_q_3_p2) - (reg_cons_4_m2 * reg_q_3_m2)) + ((reg_q_4_p2 * reg_q_3_p2) - (reg_q_4_m2 * reg_q_3_m2))))) + (0.038f * (((reg_cons_4_p3 * reg_q_3_p3) - (reg_cons_4_m3 * reg_q_3_m3)) + ((reg_q_4_p3 * reg_q_3_p3) - (reg_q_4_m3 * reg_q_3_m3))))) - (0.0035f * (((reg_cons_4_p4 * reg_q_3_p4) - (reg_cons_4_m4 * reg_q_3_m4)) + ((reg_q_4_p4 * reg_q_3_p4) - (reg_q_4_m4 * reg_q_3_m4))))) * dxinv2);
	flux_4[k*M*N + j*N + i] = r4;
    }
    __syncthreads ();
    //Value rotation


  if (j <= M-1 & i <= min (i0+blockdim_i-1, N-1)) {
    reg_cons_1_m4 = reg_cons_1_m3;
    reg_cons_1_m3 = reg_cons_1_m2;
    reg_cons_1_m2 = reg_cons_1_m1;
    reg_cons_1_m1 = sh_cons_1_c0[j-j0+4][i-i0];
    sh_cons_1_c0[j-j0+4][i-i0] = reg_cons_1_p1;
    reg_cons_1_p1 = reg_cons_1_p2;
    reg_cons_1_p2 = reg_cons_1_p3;
    reg_cons_1_p3 = reg_cons_1_p4;
    reg_cons_2_m4 = reg_cons_2_m3;
    reg_cons_2_m3 = reg_cons_2_m2;
    reg_cons_2_m2 = reg_cons_2_m1;
    reg_cons_2_m1 = sh_cons_2_c0[j-j0+4][i-i0];
    sh_cons_2_c0[j-j0+4][i-i0] = reg_cons_2_p1;
    reg_cons_2_p1 = reg_cons_2_p2;
    reg_cons_2_p2 = reg_cons_2_p3;
    reg_cons_2_p3 = reg_cons_2_p4;
    reg_cons_3_m4 = reg_cons_3_m3;
    reg_cons_3_m3 = reg_cons_3_m2;
    reg_cons_3_m2 = reg_cons_3_m1;
    reg_cons_3_m1 = sh_cons_3_c0[j-j0+4][i-i0];
    sh_cons_3_c0[j-j0+4][i-i0] = reg_cons_3_p1;
    reg_cons_3_p1 = reg_cons_3_p2;
    reg_cons_3_p2 = reg_cons_3_p3;
    reg_cons_3_p3 = reg_cons_3_p4;
    reg_cons_4_m4 = reg_cons_4_m3;
    reg_cons_4_m3 = reg_cons_4_m2;
    reg_cons_4_m2 = reg_cons_4_m1;
    reg_cons_4_m1 = sh_cons_4_c0[j-j0+4][i-i0];
    sh_cons_4_c0[j-j0+4][i-i0] = reg_cons_4_p1;
    reg_cons_4_p1 = reg_cons_4_p2;
    reg_cons_4_p2 = reg_cons_4_p3;
    reg_cons_4_p3 = reg_cons_4_p4;
    reg_q_3_m4 = reg_q_3_m3;
    reg_q_3_m3 = reg_q_3_m2;
    reg_q_3_m2 = reg_q_3_m1;
    reg_q_3_m1 = reg_q_3_c0;
    reg_q_3_c0 = reg_q_3_p1;
    reg_q_3_p1 = reg_q_3_p2;
    reg_q_3_p2 = reg_q_3_p3;
    reg_q_3_p3 = reg_q_3_p4;
    reg_q_4_m4 = reg_q_4_m3;
    reg_q_4_m3 = reg_q_4_m2;
    reg_q_4_m2 = reg_q_4_m1;
    reg_q_4_m1 = sh_q_4_c0[j-j0+4][i-i0];
    sh_q_4_c0[j-j0+4][i-i0] = reg_q_4_p1;
    reg_q_4_p1 = reg_q_4_p2;
    reg_q_4_p2 = reg_q_4_p3;
    reg_q_4_p3 = reg_q_4_p4;
   }
  }
}

extern "C" void host_code (double *h_flux_0, double *h_flux_1, double *h_flux_2, double *h_flux_3, double *h_flux_4, double *h_cons_1, double *h_cons_2, double *h_cons_3, double *h_cons_4, double *h_q_1, double *h_q_2, double *h_q_3, double *h_q_4, double dxinv0, double dxinv1, double dxinv2, int L, int M, int N) {
  double *flux_0;
  hipMalloc (&flux_0, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for flux_0\n");
  hipMemcpy (flux_0, h_flux_0, sizeof(double)*L*M*N, hipMemcpyHostToDevice);
  double *flux_1;
  hipMalloc (&flux_1, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for flux_1\n");
  hipMemcpy (flux_1, h_flux_1, sizeof(double)*L*M*N, hipMemcpyHostToDevice);
  double *flux_2;
  hipMalloc (&flux_2, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for flux_2\n");
  hipMemcpy (flux_2, h_flux_2, sizeof(double)*L*M*N, hipMemcpyHostToDevice);
  double *flux_3;
  hipMalloc (&flux_3, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for flux_3\n");
  hipMemcpy (flux_3, h_flux_3, sizeof(double)*L*M*N, hipMemcpyHostToDevice);
  double *flux_4;
  hipMalloc (&flux_4, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for flux_4\n");
  hipMemcpy (flux_4, h_flux_4, sizeof(double)*L*M*N, hipMemcpyHostToDevice);
  double *cons_1;
  hipMalloc (&cons_1, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for cons_1\n");
  hipMemcpy (cons_1, h_cons_1, sizeof(double)*L*M*N, hipMemcpyHostToDevice);
  double *cons_2;
  hipMalloc (&cons_2, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for cons_2\n");
  hipMemcpy (cons_2, h_cons_2, sizeof(double)*L*M*N, hipMemcpyHostToDevice);
  double *cons_3;
  hipMalloc (&cons_3, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for cons_3\n");
  hipMemcpy (cons_3, h_cons_3, sizeof(double)*L*M*N, hipMemcpyHostToDevice);
  double *cons_4;
  hipMalloc (&cons_4, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for cons_4\n");
  hipMemcpy (cons_4, h_cons_4, sizeof(double)*L*M*N, hipMemcpyHostToDevice);
  double *q_1;
  hipMalloc (&q_1, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for q_1\n");
  hipMemcpy (q_1, h_q_1, sizeof(double)*L*M*N, hipMemcpyHostToDevice);
  double *q_2;
  hipMalloc (&q_2, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for q_2\n");
  hipMemcpy (q_2, h_q_2, sizeof(double)*L*M*N, hipMemcpyHostToDevice);
  double *q_3;
  hipMalloc (&q_3, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for q_3\n");
  hipMemcpy (q_3, h_q_3, sizeof(double)*L*M*N, hipMemcpyHostToDevice);
  double *q_4;
  hipMalloc (&q_4, sizeof(double)*L*M*N);
  check_error ("Failed to allocate device memory for q_4\n");
  hipMemcpy (q_4, h_q_4, sizeof(double)*L*M*N, hipMemcpyHostToDevice);

  dim3 blockconfig_1 (16, 16, 1);
  dim3 gridconfig_1 (ceil(N, blockconfig_1.x-8), ceil(M-8, blockconfig_1.y), 1);
  hypterm <<<gridconfig_1, blockconfig_1>>> (flux_0, flux_1, flux_2, flux_3, flux_4, cons_1, cons_2, cons_3, cons_4, q_1, q_2, q_3, q_4, dxinv0, dxinv1, dxinv2, L, M, N);

  hipMemcpy (h_flux_0, flux_0, sizeof(double)*L*M*N, hipMemcpyDeviceToHost);
  hipMemcpy (h_flux_1, flux_1, sizeof(double)*L*M*N, hipMemcpyDeviceToHost);
  hipMemcpy (h_flux_2, flux_2, sizeof(double)*L*M*N, hipMemcpyDeviceToHost);
  hipMemcpy (h_flux_3, flux_3, sizeof(double)*L*M*N, hipMemcpyDeviceToHost);
  hipMemcpy (h_flux_4, flux_4, sizeof(double)*L*M*N, hipMemcpyDeviceToHost);

  //Free allocated memory
  hipFree (flux_0);
  hipFree (flux_1);
  hipFree (flux_2);
  hipFree (flux_3);
  hipFree (flux_4);
  hipFree (cons_1);
  hipFree (cons_2);
  hipFree (cons_3);
  hipFree (cons_4);
  hipFree (q_1);
  hipFree (q_2);
  hipFree (q_3);
  hipFree (q_4);
}

	.text
	.file	"hypterm-halo.hip"
	.globl	_Z11check_errorPKc              # -- Begin function _Z11check_errorPKc
	.type	_Z11check_errorPKc,@function
_Z11check_errorPKc:                     # @_Z11check_errorPKc
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movq	%rdi, %rbx
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB0_2
# %bb.1:
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.LBB0_2:
	.cfi_def_cfa_offset 16
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str, %edi
	movq	%rbx, %rsi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	printf
	movl	$-1, %edi
	callq	exit
.Lfunc_end0:
	.size	_Z11check_errorPKc, .Lfunc_end0-_Z11check_errorPKc
	.cfi_endproc
                                        # -- End function
	.globl	_Z22__device_stub__hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii # -- Begin function _Z22__device_stub__hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii
	.type	_Z22__device_stub__hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii,@function
_Z22__device_stub__hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii: # @_Z22__device_stub__hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii
	.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	$288, %rsp                      # imm = 0x120
	.cfi_def_cfa_offset 336
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	88(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	80(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	72(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	64(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	56(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	48(%rsp), %r8
	movq	%r9, (%r8)
	leaq	40(%rsp), %r9
	movsd	%xmm0, (%r9)
	leaq	32(%rsp), %r10
	movsd	%xmm1, (%r10)
	leaq	24(%rsp), %r11
	movsd	%xmm2, (%r11)
	leaq	128(%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)
	leaq	336(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	344(%rsp), %rax
	movq	%rax, 56(%rbx)
	leaq	352(%rsp), %rax
	movq	%rax, 64(%rbx)
	leaq	360(%rsp), %rax
	movq	%rax, 72(%rbx)
	leaq	368(%rsp), %rax
	movq	%rax, 80(%rbx)
	leaq	376(%rsp), %rax
	movq	%rax, 88(%rbx)
	leaq	384(%rsp), %rax
	movq	%rax, 96(%rbx)
	movq	%r9, 104(%rbx)
	movq	%r10, 112(%rbx)
	movq	%r11, 120(%rbx)
	leaq	392(%rsp), %rax
	movq	%rax, 128(%rbx)
	leaq	400(%rsp), %rax
	movq	%rax, 136(%rbx)
	leaq	408(%rsp), %rax
	movq	%rax, 144(%rbx)
	leaq	112(%rsp), %r14
	leaq	96(%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	$_Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$304, %rsp                      # imm = 0x130
	.cfi_adjust_cfa_offset -304
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z22__device_stub__hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii, .Lfunc_end1-_Z22__device_stub__hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii
	.cfi_endproc
                                        # -- End function
	.globl	host_code                       # -- Begin function host_code
	.type	host_code,@function
host_code:                              # @host_code
	.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	$200, %rsp
	.cfi_def_cfa_offset 256
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movsd	%xmm2, 136(%rsp)                # 8-byte Spill
	movsd	%xmm1, 120(%rsp)                # 8-byte Spill
	movsd	%xmm0, 112(%rsp)                # 8-byte Spill
	movq	%r9, 152(%rsp)                  # 8-byte Spill
	movq	%r8, 40(%rsp)                   # 8-byte Spill
	movq	%rcx, %r12
	movq	%rdx, %rbx
	movq	%rsi, %r13
	movq	%rdi, %rbp
	movslq	312(%rsp), %rax
	movslq	320(%rsp), %r15
	movq	%r15, 160(%rsp)                 # 8-byte Spill
	movslq	328(%rsp), %rcx
	movq	%rcx, 144(%rsp)                 # 8-byte Spill
	movq	%rax, 128(%rsp)                 # 8-byte Spill
	imulq	%rax, %r15
	imulq	%rcx, %r15
	shlq	$3, %r15
	leaq	32(%rsp), %r14
	movq	%r14, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movl	$.L.str.1, %edi
	callq	_Z11check_errorPKc
	movq	(%r14), %rdi
	movq	%rbp, 168(%rsp)                 # 8-byte Spill
	movq	%rbp, %rsi
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	24(%rsp), %r14
	movq	%r14, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movl	$.L.str.2, %edi
	callq	_Z11check_errorPKc
	movq	(%r14), %rdi
	movq	%r13, 176(%rsp)                 # 8-byte Spill
	movq	%r13, %rsi
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	16(%rsp), %r14
	movq	%r14, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movl	$.L.str.3, %edi
	callq	_Z11check_errorPKc
	movq	(%r14), %rdi
	movq	%rbx, 184(%rsp)                 # 8-byte Spill
	movq	%rbx, %rsi
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	8(%rsp), %r14
	movq	%r14, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movl	$.L.str.4, %edi
	callq	_Z11check_errorPKc
	movq	(%r14), %rdi
	movq	%r12, 192(%rsp)                 # 8-byte Spill
	movq	%r12, %rsi
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	%rsp, %r14
	movq	%r14, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movl	$.L.str.5, %edi
	callq	_Z11check_errorPKc
	movq	(%r14), %rdi
	movq	40(%rsp), %rsi                  # 8-byte Reload
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	104(%rsp), %r14
	movq	%r14, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movl	$.L.str.6, %edi
	callq	_Z11check_errorPKc
	movq	(%r14), %rdi
	movq	144(%rsp), %r14                 # 8-byte Reload
	movq	152(%rsp), %rsi                 # 8-byte Reload
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	96(%rsp), %rbx
	movq	%rbx, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movl	$.L.str.7, %edi
	callq	_Z11check_errorPKc
	movq	(%rbx), %rdi
	movq	256(%rsp), %rsi
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	88(%rsp), %rbx
	movq	%rbx, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movl	$.L.str.8, %edi
	callq	_Z11check_errorPKc
	movq	(%rbx), %rdi
	movq	264(%rsp), %rsi
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	80(%rsp), %rbx
	movq	%rbx, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movl	$.L.str.9, %edi
	callq	_Z11check_errorPKc
	movq	(%rbx), %rdi
	movq	272(%rsp), %rsi
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	72(%rsp), %rbx
	movq	%rbx, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movl	$.L.str.10, %edi
	callq	_Z11check_errorPKc
	movq	(%rbx), %rdi
	movq	280(%rsp), %rsi
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	64(%rsp), %rbx
	movq	%rbx, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movl	$.L.str.11, %edi
	callq	_Z11check_errorPKc
	movq	(%rbx), %rdi
	movq	288(%rsp), %rsi
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	56(%rsp), %rbx
	movq	%rbx, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movl	$.L.str.12, %edi
	callq	_Z11check_errorPKc
	movq	(%rbx), %rdi
	movq	296(%rsp), %rsi
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	48(%rsp), %rbx
	movq	%rbx, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movl	$.L.str.13, %edi
	callq	_Z11check_errorPKc
	movq	(%rbx), %rdi
	movq	160(%rsp), %rbx                 # 8-byte Reload
	movq	304(%rsp), %rsi
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movl	%r14d, %eax
	andl	$7, %eax
	movl	%r14d, %ecx
	shrl	$3, %ecx
	cmpl	$1, %eax
	sbbl	$-1, %ecx
	leal	-8(%rbx), %edi
	movl	%edi, %eax
	andl	$15, %eax
	shrl	$4, %edi
	cmpl	$1, %eax
	sbbl	$-1, %edi
	shlq	$32, %rdi
	orq	%rcx, %rdi
	movabsq	$68719476752, %rdx              # imm = 0x1000000010
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_2
# %bb.1:
	movq	32(%rsp), %rdi
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdx
	movq	8(%rsp), %rcx
	movq	(%rsp), %r8
	movq	104(%rsp), %r9
	movsd	112(%rsp), %xmm0                # 8-byte Reload
                                        # xmm0 = mem[0],zero
	movsd	120(%rsp), %xmm1                # 8-byte Reload
                                        # xmm1 = mem[0],zero
	movsd	136(%rsp), %xmm2                # 8-byte Reload
                                        # xmm2 = mem[0],zero
	pushq	%r14
	.cfi_adjust_cfa_offset 8
	pushq	%rbx
	.cfi_adjust_cfa_offset 8
	pushq	144(%rsp)                       # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	pushq	72(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	88(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	104(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	120(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	136(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	152(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	168(%rsp)
	.cfi_adjust_cfa_offset 8
	callq	_Z22__device_stub__hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii
	addq	$80, %rsp
	.cfi_adjust_cfa_offset -80
.LBB2_2:
	movq	32(%rsp), %rsi
	movq	168(%rsp), %rdi                 # 8-byte Reload
	movq	%r15, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	24(%rsp), %rsi
	movq	176(%rsp), %rdi                 # 8-byte Reload
	movq	%r15, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	16(%rsp), %rsi
	movq	184(%rsp), %rdi                 # 8-byte Reload
	movq	%r15, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	8(%rsp), %rsi
	movq	192(%rsp), %rdi                 # 8-byte Reload
	movq	%r15, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	(%rsp), %rsi
	movq	40(%rsp), %rdi                  # 8-byte Reload
	movq	%r15, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	(%rsp), %rdi
	callq	hipFree
	movq	104(%rsp), %rdi
	callq	hipFree
	movq	96(%rsp), %rdi
	callq	hipFree
	movq	88(%rsp), %rdi
	callq	hipFree
	movq	80(%rsp), %rdi
	callq	hipFree
	movq	72(%rsp), %rdi
	callq	hipFree
	movq	64(%rsp), %rdi
	callq	hipFree
	movq	56(%rsp), %rdi
	callq	hipFree
	movq	48(%rsp), %rdi
	callq	hipFree
	addq	$200, %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_end2:
	.size	host_code, .Lfunc_end2-host_code
	.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	.LBB3_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB3_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii, %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_end3:
	.size	__hip_module_ctor, .Lfunc_end3-__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	.LBB4_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
.LBB4_2:
	retq
.Lfunc_end4:
	.size	__hip_module_dtor, .Lfunc_end4-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"CUDA error : %s, %s\n"
	.size	.L.str, 21

	.type	_Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii,@object # @_Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii
	.section	.rodata,"a",@progbits
	.globl	_Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii
	.p2align	3, 0x0
_Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii:
	.quad	_Z22__device_stub__hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii
	.size	_Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii, 8

	.type	.L.str.1,@object                # @.str.1
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.1:
	.asciz	"Failed to allocate device memory for flux_0\n"
	.size	.L.str.1, 45

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"Failed to allocate device memory for flux_1\n"
	.size	.L.str.2, 45

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"Failed to allocate device memory for flux_2\n"
	.size	.L.str.3, 45

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"Failed to allocate device memory for flux_3\n"
	.size	.L.str.4, 45

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"Failed to allocate device memory for flux_4\n"
	.size	.L.str.5, 45

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"Failed to allocate device memory for cons_1\n"
	.size	.L.str.6, 45

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"Failed to allocate device memory for cons_2\n"
	.size	.L.str.7, 45

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"Failed to allocate device memory for cons_3\n"
	.size	.L.str.8, 45

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"Failed to allocate device memory for cons_4\n"
	.size	.L.str.9, 45

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"Failed to allocate device memory for q_1\n"
	.size	.L.str.10, 42

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"Failed to allocate device memory for q_2\n"
	.size	.L.str.11, 42

	.type	.L.str.12,@object               # @.str.12
.L.str.12:
	.asciz	"Failed to allocate device memory for q_3\n"
	.size	.L.str.12, 42

	.type	.L.str.13,@object               # @.str.13
.L.str.13:
	.asciz	"Failed to allocate device memory for q_4\n"
	.size	.L.str.13, 42

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii"
	.size	.L__unnamed_1, 43

	.type	__hip_fatbin_wrapper,@object    # @__hip_fatbin_wrapper
	.section	.hipFatBinSegment,"a",@progbits
	.p2align	3, 0x0
__hip_fatbin_wrapper:
	.long	1212764230                      # 0x48495046
	.long	1                               # 0x1
	.quad	__hip_fatbin
	.quad	0
	.size	__hip_fatbin_wrapper, 24

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym _Z22__device_stub__hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii ; -- Begin function _Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii
	.globl	_Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii
	.p2align	8
	.type	_Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii,@function
_Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii: ; @_Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii
; %bb.0:
	s_clause 0x3
	s_load_b32 s2, s[0:1], 0x9c
	s_load_b128 s[8:11], s[0:1], 0x80
	s_load_b512 s[16:31], s[0:1], 0x0
	s_load_b512 s[36:51], s[0:1], 0x40
	v_dual_mov_b32 v84, 0 :: v_dual_and_b32 v1, 0x3ff, v0
	v_bfe_u32 v3, v0, 10, 10
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s0, s2, 16
	s_and_b32 s2, s2, 0xffff
	s_mul_i32 s3, s15, s0
	s_add_i32 s11, s2, -8
	s_add_i32 s1, s3, 4
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_dual_mov_b32 v85, 0 :: v_dual_add_nc_u32 v2, s1, v3
	s_mul_i32 s11, s11, s14
	v_dual_mov_b32 v71, v84 :: v_dual_lshlrev_b32 v4, 3, v1
	v_dual_mov_b32 v155, v85 :: v_dual_add_nc_u32 v0, s11, v1
	v_dual_mov_b32 v154, v84 :: v_dual_mov_b32 v111, v85
	v_dual_mov_b32 v110, v84 :: v_dual_mov_b32 v113, v85
	v_dual_mov_b32 v112, v84 :: v_dual_mov_b32 v115, v85
	v_dual_mov_b32 v114, v84 :: v_dual_mov_b32 v105, v85
	v_dual_mov_b32 v104, v84 :: v_dual_mov_b32 v175, v85
	v_dual_mov_b32 v174, v84 :: v_dual_mov_b32 v179, v85
	v_dual_mov_b32 v178, v84 :: v_dual_mov_b32 v119, v85
	v_dual_mov_b32 v118, v84 :: v_dual_mov_b32 v123, v85
	v_dual_mov_b32 v122, v84 :: v_dual_mov_b32 v153, v85
	v_dual_mov_b32 v152, v84 :: v_dual_mov_b32 v135, v85
	v_dual_mov_b32 v134, v84 :: v_dual_mov_b32 v139, v85
	v_dual_mov_b32 v138, v84 :: v_dual_mov_b32 v141, v85
	v_dual_mov_b32 v140, v84 :: v_dual_mov_b32 v91, v85
	v_dual_mov_b32 v90, v84 :: v_dual_mov_b32 v145, v85
	v_dual_mov_b32 v144, v84 :: v_dual_mov_b32 v147, v85
	v_dual_mov_b32 v146, v84 :: v_dual_mov_b32 v149, v85
	v_dual_mov_b32 v148, v84 :: v_dual_mov_b32 v133, v85
	v_dual_mov_b32 v132, v84 :: v_dual_mov_b32 v103, v85
	v_dual_mov_b32 v102, v84 :: v_dual_mov_b32 v107, v85
	v_dual_mov_b32 v106, v84 :: v_dual_mov_b32 v109, v85
	v_dual_mov_b32 v108, v84 :: v_dual_mov_b32 v129, v85
	v_dual_mov_b32 v128, v84 :: v_dual_mov_b32 v143, v85
	v_dual_mov_b32 v142, v84 :: v_dual_mov_b32 v191, v85
	s_add_i32 s0, s11, s2
	v_dual_mov_b32 v190, v84 :: v_dual_mov_b32 v93, v85
	v_dual_mov_b32 v92, v84 :: v_dual_mov_b32 v99, v85
	s_min_i32 s5, s0, s10
	v_dual_mov_b32 v98, v84 :: v_dual_mov_b32 v127, v85
	v_dual_mov_b32 v126, v84 :: v_dual_mov_b32 v125, v85
	v_cmp_gt_i32_e32 vcc_lo, s9, v2
	v_cmp_gt_i32_e64 s2, s5, v0
	v_dual_mov_b32 v124, v84 :: v_dual_mov_b32 v121, v85
	v_dual_mov_b32 v72, v85 :: v_dual_mov_b32 v73, v84
	v_dual_mov_b32 v120, v84 :: v_dual_mov_b32 v171, v85
	v_dual_mov_b32 v76, v84 :: v_dual_mov_b32 v77, v85
	v_dual_mov_b32 v188, v85 :: v_dual_mov_b32 v187, v84
	v_dual_mov_b32 v157, v85 :: v_dual_mov_b32 v156, v84
	v_dual_mov_b32 v87, v85 :: v_dual_mov_b32 v86, v84
	v_dual_mov_b32 v95, v85 :: v_dual_mov_b32 v94, v84
	v_dual_mov_b32 v97, v85 :: v_dual_mov_b32 v96, v84
	v_dual_mov_b32 v101, v85 :: v_dual_mov_b32 v100, v84
	v_dual_mov_b32 v82, v84 :: v_dual_mov_b32 v83, v85
	v_dual_mov_b32 v78, v84 :: v_dual_mov_b32 v79, v85
	v_dual_mov_b32 v80, v84 :: v_dual_mov_b32 v81, v85
	v_mov_b32_e32 v74, v85
	v_mov_b32_e32 v170, v84
	s_and_b32 s14, vcc_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s0, s14
	s_cbranch_execz .LBB0_2
; %bb.1:
	v_add_nc_u32_e32 v11, s9, v2
	v_mad_u64_u32 v[7:8], null, v2, s10, v[0:1]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v13, s9, v11
	v_mad_u64_u32 v[9:10], null, v11, s10, v[0:1]
	v_ashrrev_i32_e32 v8, 31, v7
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v15, s9, v13
	v_mad_u64_u32 v[11:12], null, v13, s10, v[0:1]
	v_lshlrev_b64 v[7:8], 3, v[7:8]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_3) | instid1(VALU_DEP_4)
	v_add_nc_u32_e32 v35, s9, v15
	v_mad_u64_u32 v[13:14], null, v15, s10, v[0:1]
	v_ashrrev_i32_e32 v10, 31, v9
	v_ashrrev_i32_e32 v12, 31, v11
	v_mad_u64_u32 v[5:6], null, v35, s10, v[0:1]
	v_add_nc_u32_e32 v47, s9, v35
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_lshlrev_b64 v[9:10], 3, v[9:10]
	v_lshlrev_b64 v[11:12], 3, v[11:12]
	v_ashrrev_i32_e32 v14, 31, v13
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_mad_u64_u32 v[45:46], null, v47, s10, v[0:1]
	v_ashrrev_i32_e32 v6, 31, v5
	v_add_nc_u32_e32 v46, s9, v47
	v_lshlrev_b64 v[13:14], 3, v[13:14]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[5:6], 3, v[5:6]
	v_mad_u64_u32 v[57:58], null, v46, s10, v[0:1]
	v_add_nc_u32_e32 v58, s9, v46
	v_ashrrev_i32_e32 v46, 31, v45
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_u32 v15, vcc_lo, s26, v5
	v_add_co_ci_u32_e32 v16, vcc_lo, s27, v6, vcc_lo
	v_add_co_u32 v17, vcc_lo, s28, v5
	v_add_co_ci_u32_e32 v18, vcc_lo, s29, v6, vcc_lo
	v_add_co_u32 v19, vcc_lo, s30, v5
	v_add_co_ci_u32_e32 v20, vcc_lo, s31, v6, vcc_lo
	v_add_co_u32 v21, vcc_lo, s36, v5
	v_add_co_ci_u32_e32 v22, vcc_lo, s37, v6, vcc_lo
	v_add_co_u32 v23, vcc_lo, s44, v5
	v_add_co_ci_u32_e32 v24, vcc_lo, s45, v6, vcc_lo
	v_add_co_u32 v25, vcc_lo, s26, v7
	v_add_co_ci_u32_e32 v26, vcc_lo, s27, v8, vcc_lo
	v_add_co_u32 v27, vcc_lo, s28, v7
	v_add_co_ci_u32_e32 v28, vcc_lo, s29, v8, vcc_lo
	v_add_co_u32 v29, vcc_lo, s30, v7
	v_add_co_ci_u32_e32 v30, vcc_lo, s31, v8, vcc_lo
	v_add_co_u32 v31, vcc_lo, s36, v7
	v_add_co_ci_u32_e32 v32, vcc_lo, s37, v8, vcc_lo
	v_add_co_u32 v33, vcc_lo, s42, v7
	v_add_co_ci_u32_e32 v34, vcc_lo, s43, v8, vcc_lo
	v_add_co_u32 v7, vcc_lo, s44, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s45, v8, vcc_lo
	v_add_co_u32 v35, vcc_lo, s26, v9
	v_add_co_ci_u32_e32 v36, vcc_lo, s27, v10, vcc_lo
	v_add_co_u32 v37, vcc_lo, s28, v9
	v_add_co_ci_u32_e32 v38, vcc_lo, s29, v10, vcc_lo
	v_add_co_u32 v39, vcc_lo, s30, v9
	v_add_co_ci_u32_e32 v40, vcc_lo, s31, v10, vcc_lo
	v_add_co_u32 v41, vcc_lo, s36, v9
	v_add_co_ci_u32_e32 v42, vcc_lo, s37, v10, vcc_lo
	v_add_co_u32 v43, vcc_lo, s42, v9
	v_add_co_ci_u32_e32 v44, vcc_lo, s43, v10, vcc_lo
	v_add_co_u32 v9, vcc_lo, s44, v9
	v_add_co_ci_u32_e32 v10, vcc_lo, s45, v10, vcc_lo
	v_add_co_u32 v47, vcc_lo, s26, v11
	v_add_co_ci_u32_e32 v48, vcc_lo, s27, v12, vcc_lo
	v_add_co_u32 v49, vcc_lo, s28, v11
	v_add_co_ci_u32_e32 v50, vcc_lo, s29, v12, vcc_lo
	v_add_co_u32 v51, vcc_lo, s30, v11
	v_add_co_ci_u32_e32 v52, vcc_lo, s31, v12, vcc_lo
	v_add_co_u32 v53, vcc_lo, s36, v11
	v_add_co_ci_u32_e32 v54, vcc_lo, s37, v12, vcc_lo
	v_add_co_u32 v55, vcc_lo, s42, v11
	v_add_co_ci_u32_e32 v56, vcc_lo, s43, v12, vcc_lo
	v_add_co_u32 v11, vcc_lo, s44, v11
	v_add_co_ci_u32_e32 v12, vcc_lo, s45, v12, vcc_lo
	v_add_co_u32 v59, vcc_lo, s26, v13
	v_add_co_ci_u32_e32 v60, vcc_lo, s27, v14, vcc_lo
	v_add_co_u32 v61, vcc_lo, s28, v13
	v_lshlrev_b64 v[45:46], 3, v[45:46]
	v_add_co_ci_u32_e32 v62, vcc_lo, s29, v14, vcc_lo
	v_add_co_u32 v65, vcc_lo, s30, v13
	v_add_co_ci_u32_e32 v66, vcc_lo, s31, v14, vcc_lo
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_u32 v67, vcc_lo, s26, v45
	v_add_co_ci_u32_e32 v68, vcc_lo, s27, v46, vcc_lo
	v_add_co_u32 v69, vcc_lo, s28, v45
	v_mad_u64_u32 v[63:64], null, v58, s10, v[0:1]
	v_ashrrev_i32_e32 v58, 31, v57
	v_add_co_ci_u32_e32 v70, vcc_lo, s29, v46, vcc_lo
	v_add_co_u32 v71, vcc_lo, s30, v45
	v_add_co_ci_u32_e32 v72, vcc_lo, s31, v46, vcc_lo
	s_delay_alu instid0(VALU_DEP_4)
	v_lshlrev_b64 v[57:58], 3, v[57:58]
	v_add_co_u32 v73, vcc_lo, s36, v45
	v_add_co_ci_u32_e32 v74, vcc_lo, s37, v46, vcc_lo
	v_add_co_u32 v84, vcc_lo, s44, v45
	v_add_co_ci_u32_e32 v85, vcc_lo, s45, v46, vcc_lo
	v_add_co_u32 v75, vcc_lo, s26, v57
	v_add_co_ci_u32_e32 v76, vcc_lo, s27, v58, vcc_lo
	v_add_co_u32 v78, vcc_lo, s28, v57
	v_ashrrev_i32_e32 v64, 31, v63
	v_add_co_ci_u32_e32 v79, vcc_lo, s29, v58, vcc_lo
	v_add_co_u32 v80, vcc_lo, s30, v57
	v_add_co_ci_u32_e32 v81, vcc_lo, s31, v58, vcc_lo
	s_delay_alu instid0(VALU_DEP_4)
	v_lshlrev_b64 v[63:64], 3, v[63:64]
	v_add_co_u32 v86, vcc_lo, s36, v57
	v_add_co_ci_u32_e32 v87, vcc_lo, s37, v58, vcc_lo
	v_add_co_u32 v94, vcc_lo, s44, v57
	v_add_co_ci_u32_e32 v95, vcc_lo, s45, v58, vcc_lo
	v_add_co_u32 v82, vcc_lo, s26, v63
	v_add_co_ci_u32_e32 v83, vcc_lo, s27, v64, vcc_lo
	v_add_co_u32 v88, vcc_lo, s28, v63
	v_add_co_ci_u32_e32 v89, vcc_lo, s29, v64, vcc_lo
	s_clause 0x3
	global_load_b64 v[15:16], v[15:16], off
	global_load_b64 v[98:99], v[67:68], off
	global_load_b64 v[92:93], v[75:76], off
	global_load_b64 v[190:191], v[82:83], off
	s_clause 0x3
	global_load_b64 v[17:18], v[17:18], off
	global_load_b64 v[76:77], v[69:70], off
	global_load_b64 v[187:188], v[78:79], off
	global_load_b64 v[156:157], v[88:89], off
	v_add_co_u32 v67, vcc_lo, s30, v63
	v_add_co_ci_u32_e32 v68, vcc_lo, s31, v64, vcc_lo
	v_add_co_u32 v69, vcc_lo, s36, v63
	v_add_co_ci_u32_e32 v70, vcc_lo, s37, v64, vcc_lo
	s_clause 0x3
	global_load_b64 v[19:20], v[19:20], off
	global_load_b64 v[82:83], v[71:72], off
	global_load_b64 v[78:79], v[80:81], off
	global_load_b64 v[80:81], v[67:68], off
	s_clause 0x3
	global_load_b64 v[21:22], v[21:22], off
	global_load_b64 v[104:105], v[73:74], off
	global_load_b64 v[71:72], v[86:87], off
	global_load_b64 v[73:74], v[69:70], off
	v_add_co_u32 v67, vcc_lo, s44, v63
	v_add_co_ci_u32_e32 v68, vcc_lo, s45, v64, vcc_lo
	v_add_co_u32 v69, vcc_lo, s36, v13
	v_add_co_ci_u32_e32 v70, vcc_lo, s37, v14, vcc_lo
	s_clause 0x3
	global_load_b64 v[23:24], v[23:24], off
	global_load_b64 v[132:133], v[84:85], off
	global_load_b64 v[102:103], v[94:95], off
	global_load_b64 v[106:107], v[67:68], off
	v_add_co_u32 v67, vcc_lo, s42, v13
	v_add_co_ci_u32_e32 v68, vcc_lo, s43, v14, vcc_lo
	v_add_co_u32 v13, vcc_lo, s44, v13
	v_add_co_ci_u32_e32 v14, vcc_lo, s45, v14, vcc_lo
	v_add_co_u32 v5, vcc_lo, s42, v5
	v_add_co_ci_u32_e32 v6, vcc_lo, s43, v6, vcc_lo
	v_add_co_u32 v45, vcc_lo, s42, v45
	v_add_co_ci_u32_e32 v46, vcc_lo, s43, v46, vcc_lo
	v_add_co_u32 v57, vcc_lo, s42, v57
	v_add_co_ci_u32_e32 v58, vcc_lo, s43, v58, vcc_lo
	v_add_co_u32 v63, vcc_lo, s42, v63
	v_add_co_ci_u32_e32 v64, vcc_lo, s43, v64, vcc_lo
	s_clause 0x3
	global_load_b64 v[170:171], v[25:26], off
	global_load_b64 v[120:121], v[35:36], off
	global_load_b64 v[124:125], v[47:48], off
	global_load_b64 v[126:127], v[59:60], off
	s_clause 0x3
	global_load_b64 v[142:143], v[27:28], off
	global_load_b64 v[128:129], v[37:38], off
	global_load_b64 v[108:109], v[49:50], off
	global_load_b64 v[84:85], v[61:62], off
	s_clause 0x3
	global_load_b64 v[86:87], v[29:30], off
	global_load_b64 v[94:95], v[39:40], off
	global_load_b64 v[96:97], v[51:52], off
	global_load_b64 v[100:101], v[65:66], off
	s_clause 0x3
	global_load_b64 v[154:155], v[31:32], off
	global_load_b64 v[110:111], v[41:42], off
	global_load_b64 v[112:113], v[53:54], off
	global_load_b64 v[114:115], v[69:70], off
	s_clause 0x7
	global_load_b64 v[174:175], v[33:34], off
	global_load_b64 v[178:179], v[43:44], off
	global_load_b64 v[118:119], v[55:56], off
	global_load_b64 v[122:123], v[67:68], off
	global_load_b64 v[152:153], v[5:6], off
	global_load_b64 v[134:135], v[45:46], off
	global_load_b64 v[138:139], v[57:58], off
	global_load_b64 v[140:141], v[63:64], off
	s_clause 0x3
	global_load_b64 v[90:91], v[7:8], off
	global_load_b64 v[144:145], v[9:10], off
	global_load_b64 v[146:147], v[11:12], off
	global_load_b64 v[148:149], v[13:14], off
	v_lshl_add_u32 v5, v3, 7, v4
	s_waitcnt vmcnt(43)
	ds_store_2addr_stride64_b64 v5, v[15:16], v[17:18] offset0:1 offset1:7
	s_waitcnt vmcnt(35)
	ds_store_2addr_stride64_b64 v5, v[19:20], v[21:22] offset0:13 offset1:19
	s_waitcnt vmcnt(31)
	ds_store_b64 v5, v[23:24] offset:18944
.LBB0_2:
	s_or_b32 exec_lo, exec_lo, s0
	v_cmp_lt_u32_e32 vcc_lo, 3, v3
	s_mov_b32 s4, 0
                                        ; implicit-def: $vgpr6
	s_and_saveexec_b32 s0, vcc_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s6, exec_lo, s0
	s_cbranch_execz .LBB0_6
; %bb.3:
	s_mov_b32 s7, exec_lo
                                        ; implicit-def: $vgpr6
	v_cmpx_gt_u32_e32 8, v3
	s_xor_b32 s7, exec_lo, s7
; %bb.4:
	v_add3_u32 v6, s3, v3, 16
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_gt_i32_e64 s0, s9, v6
	s_and_b32 s0, s0, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 s4, s0, exec_lo
; %bb.5:                                ; %Flow2212
	s_or_b32 exec_lo, exec_lo, s7
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 s4, s4, exec_lo
.LBB0_6:                                ; %Flow2211
	s_or_saveexec_b32 s6, s6
	v_add_nc_u32_e32 v5, s3, v3
	s_xor_b32 exec_lo, exec_lo, s6
; %bb.7:
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_cmp_gt_i32_e64 s0, s9, v5
	v_mov_b32_e32 v6, v5
	s_and_not1_b32 s3, s4, exec_lo
	s_and_b32 s0, s0, s2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, s0, exec_lo
	s_or_b32 s4, s3, s0
; %bb.8:                                ; %Flow2213
	s_or_b32 exec_lo, exec_lo, s6
	s_and_saveexec_b32 s3, s4
	s_cbranch_execz .LBB0_10
; %bb.9:                                ; %.sink.split
	v_lshl_add_u32 v9, s9, 2, v6
	v_subrev_nc_u32_e32 v6, s1, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mad_u64_u32 v[7:8], null, v9, s10, v[0:1]
	v_lshl_add_u32 v6, v6, 7, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v8, 31, v7
	v_lshlrev_b64 v[7:8], 3, v[7:8]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v9, s0, s26, v7
	v_add_co_ci_u32_e64 v10, s0, s27, v8, s0
	v_add_co_u32 v11, s0, s28, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e64 v12, s0, s29, v8, s0
	v_add_co_u32 v13, s0, s30, v7
	v_add_co_ci_u32_e64 v14, s0, s31, v8, s0
	v_add_co_u32 v15, s0, s36, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e64 v16, s0, s37, v8, s0
	v_add_co_u32 v7, s0, s44, v7
	v_add_co_ci_u32_e64 v8, s0, s45, v8, s0
	global_load_b64 v[9:10], v[9:10], off
	global_load_b64 v[11:12], v[11:12], off
	global_load_b64 v[13:14], v[13:14], off
	global_load_b64 v[15:16], v[15:16], off
	global_load_b64 v[7:8], v[7:8], off
	s_waitcnt vmcnt(3)
	ds_store_2addr_stride64_b64 v6, v[9:10], v[11:12] offset0:1 offset1:7
	s_waitcnt vmcnt(1)
	ds_store_2addr_stride64_b64 v6, v[13:14], v[15:16] offset0:13 offset1:19
	s_waitcnt vmcnt(0)
	ds_store_b64 v6, v[7:8] offset:18944
.LBB0_10:
	s_or_b32 exec_lo, exec_lo, s3
	s_cmp_lt_i32 s8, 9
	s_cbranch_scc1 .LBB0_41
; %bb.11:                               ; %.lr.ph
	v_lshl_add_u32 v168, v3, 7, v4
	v_cmp_lt_u32_e64 s0, 7, v3
	v_subrev_nc_u32_e32 v7, s1, v5
	v_cmp_lt_u32_e64 s1, 11, v3
	v_dual_mov_b32 v166, 0 :: v_dual_add_nc_u32 v3, 16, v5
	v_dual_mov_b32 v167, 0 :: v_dual_add_nc_u32 v6, 4, v2
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshl_add_u32 v4, v7, 7, v4
	v_cmp_gt_i32_e64 s4, s9, v3
	s_delay_alu instid0(VALU_DEP_4)
	v_dual_mov_b32 v88, v166 :: v_dual_add_nc_u32 v3, 0x4c00, v168
	s_max_i32 s6, s11, 0
	v_cmp_gt_i32_e64 s3, s9, v6
	v_dual_mov_b32 v89, v167 :: v_dual_add_nc_u32 v6, -12, v2
	scratch_store_b32 off, v3, off offset:64 ; 4-byte Folded Spill
	v_dual_mov_b32 v12, v166 :: v_dual_add_nc_u32 v3, 0x2c00, v4
	s_add_i32 s5, s5, -5
	s_add_i32 s6, s6, 4
	s_and_b32 s34, s4, s2
	scratch_store_b32 off, v3, off offset:72 ; 4-byte Folded Spill
	v_add_nc_u32_e32 v3, 0x4400, v168
	v_cmp_ge_i32_e64 s4, s5, v0
	v_cmp_le_i32_e64 s6, s6, v0
	v_dual_mov_b32 v13, v167 :: v_dual_add_nc_u32 v0, 0x2e00, v4
	scratch_store_b32 off, v3, off offset:68 ; 4-byte Folded Spill
	v_dual_mov_b32 v150, v166 :: v_dual_add_nc_u32 v3, 0x200, v4
	s_and_b32 s15, s2, s3
	v_cmp_gt_i32_e64 s3, s9, v6
	s_clause 0x1
	scratch_store_b32 off, v0, off offset:60
	scratch_store_b32 off, v3, off offset:44
	v_dual_mov_b32 v116, v166 :: v_dual_add_nc_u32 v3, 0xe00, v4
	v_lshl_add_u32 v0, s9, 2, v5
	s_and_b32 s33, s2, s3
	s_add_i32 s3, s9, -5
	scratch_store_b32 off, v3, off offset:48 ; 4-byte Folded Spill
	v_add_nc_u32_e32 v3, 0x1a00, v4
	v_cmp_ge_i32_e64 s3, s3, v2
	v_dual_mov_b32 v117, v167 :: v_dual_add_nc_u32 v2, 8, v0
	v_cmp_gt_i32_e64 s5, s9, v5
	scratch_store_b32 off, v3, off offset:56 ; 4-byte Folded Spill
	v_add_nc_u32_e32 v3, 0x2600, v4
	v_mad_u64_u32 v[158:159], null, s10, v2, v[1:2]
	v_add_nc_u32_e32 v2, 16, v0
	s_and_b32 s35, s5, s2
	scratch_store_b32 off, v3, off offset:52 ; 4-byte Folded Spill
	v_add_nc_u32_e32 v3, -8, v0
	s_lshl_b32 s2, s9, 3
	v_add_nc_u32_e32 v130, 0x3000, v168
	s_and_b32 s3, s3, s4
	v_add_nc_u32_e32 v189, 0x4800, v168
	v_mad_u64_u32 v[159:160], null, s10, v3, v[1:2]
	v_add_nc_u32_e32 v3, 4, v0
	v_mad_u64_u32 v[160:161], null, s10, v2, v[1:2]
	v_add3_u32 v2, v5, s2, 4
	v_dual_mov_b32 v151, v167 :: v_dual_add_nc_u32 v186, 0xc00, v168
	v_add_nc_u32_e32 v131, 0x1800, v168
	v_add_nc_u32_e32 v136, 0x2400, v168
	s_delay_alu instid0(VALU_DEP_4)
	v_mad_u64_u32 v[161:162], null, s10, v3, v[1:2]
	v_mad_u64_u32 v[162:163], null, s10, v0, v[1:2]
	v_mad_u64_u32 v[163:164], null, s10, v2, v[1:2]
	v_dual_mov_b32 v164, v166 :: v_dual_mov_b32 v165, v167
	s_and_b32 s52, s3, s6
	s_brev_b32 s4, 5
	s_brev_b32 s6, 6
	s_mov_b32 s12, 2.0
	s_mov_b32 s5, 0x3fc99999
	s_mov_b32 s7, 0x3fa374bc
	s_mov_b32 s13, 0xbf6cac08
	s_add_i32 s8, s8, -8
	s_mul_i32 s9, s10, s9
	s_mov_b32 s3, 0x3fe99999
.LBB0_12:                               ; =>This Inner Loop Header: Depth=1
	v_add_nc_u32_e32 v0, s11, v161
	s_delay_alu instid0(VALU_DEP_1)
	v_ashrrev_i32_e32 v1, 31, v0
	s_and_saveexec_b32 s10, s14
	s_cbranch_execz .LBB0_14
; %bb.13:                               ;   in Loop: Header=BB0_12 Depth=1
	v_add_nc_u32_e32 v2, s11, v163
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[10:11], 3, v[0:1]
	v_ashrrev_i32_e32 v3, 31, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[2:3], 3, v[2:3]
	v_add_co_u32 v4, s2, s26, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e64 v5, s2, s27, v3, s2
	v_add_co_u32 v6, s2, s28, v2
	v_add_co_ci_u32_e64 v7, s2, s29, v3, s2
	v_add_co_u32 v8, s2, s30, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e64 v9, s2, s31, v3, s2
	global_load_b64 v[166:167], v[4:5], off
	global_load_b64 v[116:117], v[6:7], off
	global_load_b64 v[150:151], v[8:9], off
	v_add_co_u32 v4, s2, s36, v2
	v_add_co_ci_u32_e64 v5, s2, s37, v3, s2
	v_add_co_u32 v6, s2, s38, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e64 v7, s2, s39, v11, s2
	v_add_co_u32 v8, s2, s40, v10
	v_add_co_ci_u32_e64 v9, s2, s41, v11, s2
	v_add_co_u32 v10, s2, s42, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e64 v11, s2, s43, v3, s2
	v_add_co_u32 v2, s2, s44, v2
	v_add_co_ci_u32_e64 v3, s2, s45, v3, s2
	global_load_b64 v[12:13], v[4:5], off
	global_load_b64 v[4:5], v[6:7], off
	global_load_b64 v[6:7], v[8:9], off
	global_load_b64 v[88:89], v[10:11], off
	global_load_b64 v[164:165], v[2:3], off
	v_add_nc_u32_e32 v2, 0x3e00, v168
	s_waitcnt vmcnt(3)
	ds_store_b64 v130, v[4:5] offset:512
	s_waitcnt vmcnt(2)
	ds_store_b64 v2, v[6:7]
.LBB0_14:                               ;   in Loop: Header=BB0_12 Depth=1
	s_or_b32 exec_lo, exec_lo, s10
	s_mov_b32 s53, 0
                                        ; implicit-def: $vgpr4
                                        ; implicit-def: $vgpr2_vgpr3
	s_and_saveexec_b32 s2, vcc_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s10, exec_lo, s2
	s_cbranch_execz .LBB0_30
; %bb.15:                               ;   in Loop: Header=BB0_12 Depth=1
	s_mov_b32 s54, 0
                                        ; implicit-def: $vgpr4
                                        ; implicit-def: $vgpr2_vgpr3
	s_and_saveexec_b32 s2, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s53, exec_lo, s2
	s_cbranch_execz .LBB0_25
; %bb.16:                               ;   in Loop: Header=BB0_12 Depth=1
                                        ; implicit-def: $vgpr4
                                        ; implicit-def: $vgpr2_vgpr3
	s_and_saveexec_b32 s2, s1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s55, exec_lo, s2
	s_cbranch_execz .LBB0_20
; %bb.17:                               ;   in Loop: Header=BB0_12 Depth=1
                                        ; implicit-def: $vgpr2_vgpr3
	s_and_saveexec_b32 s2, s15
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s56, exec_lo, s2
	s_cbranch_execz .LBB0_19
; %bb.18:                               ;   in Loop: Header=BB0_12 Depth=1
	v_add_nc_u32_e32 v2, s11, v158
	s_mov_b32 s54, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	v_lshlrev_b64 v[2:3], 3, v[2:3]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v4, s2, s38, v2
	v_add_co_ci_u32_e64 v5, s2, s39, v3, s2
	v_add_co_u32 v6, s2, s40, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v7, s2, s41, v3, s2
	v_add_co_u32 v2, s2, s44, v2
	global_load_b64 v[4:5], v[4:5], off
	global_load_b64 v[6:7], v[6:7], off
	v_add_co_ci_u32_e64 v3, s2, s45, v3, s2
	s_waitcnt vmcnt(1)
	ds_store_b64 v130, v[4:5] offset:1024
	v_add_nc_u32_e32 v4, 0x4000, v168
	s_waitcnt vmcnt(0)
	ds_store_b64 v4, v[6:7]
.LBB0_19:                               ; %Flow2202
                                        ;   in Loop: Header=BB0_12 Depth=1
	s_or_b32 exec_lo, exec_lo, s56
	scratch_load_b32 v4, off, off offset:64 ; 4-byte Folded Reload
	s_and_b32 s54, s54, exec_lo
.LBB0_20:                               ; %Flow2201
                                        ;   in Loop: Header=BB0_12 Depth=1
	s_and_not1_saveexec_b32 s55, s55
	s_cbranch_execz .LBB0_24
; %bb.21:                               ;   in Loop: Header=BB0_12 Depth=1
	s_mov_b32 s2, s54
                                        ; implicit-def: $vgpr2_vgpr3
	s_and_saveexec_b32 s56, s33
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s56, exec_lo, s56
	s_cbranch_execz .LBB0_23
; %bb.22:                               ;   in Loop: Header=BB0_12 Depth=1
	v_add_nc_u32_e32 v2, s11, v159
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	v_lshlrev_b64 v[2:3], 3, v[2:3]
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v4, s2, s38, v2
	v_add_co_ci_u32_e64 v5, s2, s39, v3, s2
	v_add_co_u32 v6, s2, s40, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v7, s2, s41, v3, s2
	v_add_co_u32 v2, s2, s44, v2
	global_load_b64 v[4:5], v[4:5], off
	global_load_b64 v[6:7], v[6:7], off
	scratch_load_b32 v8, off, off offset:72 ; 4-byte Folded Reload
	v_add_co_ci_u32_e64 v3, s2, s45, v3, s2
	s_or_b32 s2, s54, exec_lo
	s_waitcnt vmcnt(0)
	ds_store_b64 v8, v[4:5] offset:512
	ds_store_b64 v130, v[6:7] offset:2048
.LBB0_23:                               ; %Flow2204
                                        ;   in Loop: Header=BB0_12 Depth=1
	s_or_b32 exec_lo, exec_lo, s56
	scratch_load_b32 v4, off, off offset:68 ; 4-byte Folded Reload
	s_and_not1_b32 s54, s54, exec_lo
	s_and_b32 s2, s2, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s54, s54, s2
.LBB0_24:                               ; %Flow2203
                                        ;   in Loop: Header=BB0_12 Depth=1
	s_or_b32 exec_lo, exec_lo, s55
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 s54, s54, exec_lo
.LBB0_25:                               ; %Flow2200
                                        ;   in Loop: Header=BB0_12 Depth=1
	s_and_not1_saveexec_b32 s53, s53
	s_cbranch_execz .LBB0_29
; %bb.26:                               ;   in Loop: Header=BB0_12 Depth=1
	s_mov_b32 s2, s54
                                        ; implicit-def: $vgpr2_vgpr3
	s_and_saveexec_b32 s55, s34
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s55, exec_lo, s55
	s_cbranch_execz .LBB0_28
; %bb.27:                               ;   in Loop: Header=BB0_12 Depth=1
	v_add_nc_u32_e32 v2, s11, v160
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	v_lshlrev_b64 v[2:3], 3, v[2:3]
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v4, s2, s26, v2
	v_add_co_ci_u32_e64 v5, s2, s27, v3, s2
	v_add_co_u32 v6, s2, s28, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e64 v7, s2, s29, v3, s2
	v_add_co_u32 v8, s2, s30, v2
	v_add_co_ci_u32_e64 v9, s2, s31, v3, s2
	global_load_b64 v[4:5], v[4:5], off
	global_load_b64 v[6:7], v[6:7], off
	global_load_b64 v[8:9], v[8:9], off
	scratch_load_b32 v10, off, off offset:44 ; 4-byte Folded Reload
	v_add_co_u32 v2, s2, s36, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v3, s2, s37, v3, s2
	s_or_b32 s2, s54, exec_lo
	s_waitcnt vmcnt(0)
	ds_store_b64 v10, v[4:5] offset:2048
	scratch_load_b32 v4, off, off offset:48 ; 4-byte Folded Reload
	s_waitcnt vmcnt(0)
	ds_store_b64 v4, v[6:7] offset:2048
	scratch_load_b32 v4, off, off offset:56 ; 4-byte Folded Reload
	s_waitcnt vmcnt(0)
	ds_store_b64 v4, v[8:9] offset:2048
.LBB0_28:                               ; %Flow2206
                                        ;   in Loop: Header=BB0_12 Depth=1
	s_or_b32 exec_lo, exec_lo, s55
	scratch_load_b32 v4, off, off offset:60 ; 4-byte Folded Reload
	s_and_not1_b32 s54, s54, exec_lo
	s_and_b32 s2, s2, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s54, s54, s2
.LBB0_29:                               ; %Flow2205
                                        ;   in Loop: Header=BB0_12 Depth=1
	s_or_b32 exec_lo, exec_lo, s53
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 s53, s54, exec_lo
.LBB0_30:                               ; %Flow
                                        ;   in Loop: Header=BB0_12 Depth=1
	s_and_not1_saveexec_b32 s10, s10
	s_cbranch_execz .LBB0_34
; %bb.31:                               ;   in Loop: Header=BB0_12 Depth=1
	s_mov_b32 s2, s53
                                        ; implicit-def: $vgpr2_vgpr3
	s_and_saveexec_b32 s54, s35
	s_cbranch_execz .LBB0_33
; %bb.32:                               ;   in Loop: Header=BB0_12 Depth=1
	v_add_nc_u32_e32 v2, s11, v162
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	v_lshlrev_b64 v[2:3], 3, v[2:3]
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v4, s2, s26, v2
	v_add_co_ci_u32_e64 v5, s2, s27, v3, s2
	v_add_co_u32 v6, s2, s28, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e64 v7, s2, s29, v3, s2
	v_add_co_u32 v8, s2, s30, v2
	v_add_co_ci_u32_e64 v9, s2, s31, v3, s2
	global_load_b64 v[4:5], v[4:5], off
	global_load_b64 v[6:7], v[6:7], off
	global_load_b64 v[8:9], v[8:9], off
	scratch_load_b32 v10, off, off offset:44 ; 4-byte Folded Reload
	v_add_co_u32 v2, s2, s36, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v3, s2, s37, v3, s2
	s_or_b32 s2, s53, exec_lo
	s_waitcnt vmcnt(0)
	ds_store_b64 v10, v[4:5]
	scratch_load_b32 v4, off, off offset:48 ; 4-byte Folded Reload
	s_waitcnt vmcnt(0)
	ds_store_b64 v4, v[6:7]
	scratch_load_b32 v4, off, off offset:56 ; 4-byte Folded Reload
	s_waitcnt vmcnt(0)
	ds_store_b64 v4, v[8:9]
.LBB0_33:                               ; %Flow2208
                                        ;   in Loop: Header=BB0_12 Depth=1
	s_or_b32 exec_lo, exec_lo, s54
	scratch_load_b32 v4, off, off offset:52 ; 4-byte Folded Reload
	s_and_not1_b32 s53, s53, exec_lo
	s_and_b32 s2, s2, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s53, s53, s2
.LBB0_34:                               ; %Flow2207
                                        ;   in Loop: Header=BB0_12 Depth=1
	s_or_b32 exec_lo, exec_lo, s10
	s_and_saveexec_b32 s2, s53
	s_cbranch_execz .LBB0_36
; %bb.35:                               ; %.sink.split2019
                                        ;   in Loop: Header=BB0_12 Depth=1
	global_load_b64 v[2:3], v[2:3], off
	s_waitcnt vmcnt(0)
	ds_store_b64 v4, v[2:3]
.LBB0_36:                               ;   in Loop: Header=BB0_12 Depth=1
	s_or_b32 exec_lo, exec_lo, s2
	s_waitcnt vmcnt(0) lgkmcnt(0)
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_and_saveexec_b32 s10, s52
	s_cbranch_execz .LBB0_38
; %bb.37:                               ;   in Loop: Header=BB0_12 Depth=1
	v_lshlrev_b64 v[0:1], 3, v[0:1]
	v_add_nc_u32_e32 v137, 0x3c00, v168
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v16, s2, s16, v0
	v_add_co_ci_u32_e64 v17, s2, s17, v1, s2
	v_add_co_u32 v62, s2, s18, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e64 v63, s2, s19, v1, s2
	v_add_co_u32 v184, s2, s20, v0
	v_add_co_ci_u32_e64 v185, s2, s21, v1, s2
	v_add_co_u32 v180, s2, s22, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e64 v181, s2, s23, v1, s2
	v_add_co_u32 v176, s2, s24, v0
	v_add_co_ci_u32_e64 v177, s2, s25, v1, s2
	global_load_b64 v[2:3], v[180:181], off
	global_load_b64 v[18:19], v[16:17], off
	global_load_b64 v[58:59], v[62:63], off
	global_load_b64 v[0:1], v[176:177], off
	global_load_b64 v[56:57], v[184:185], off
	s_mov_b32 s2, s4
	s_waitcnt vmcnt(4)
	scratch_store_b64 off, v[2:3], off offset:36 ; 8-byte Folded Spill
	s_waitcnt vmcnt(1)
	s_clause 0x1
	scratch_store_b64 off, v[0:1], off offset:28
	scratch_store_b64 off, v[12:13], off offset:20
	scratch_load_b32 v12, off, off offset:44 ; 4-byte Folded Reload
	s_waitcnt vmcnt(0)
	ds_load_2addr_b64 v[0:3], v12 offset0:65 offset1:66
	ds_load_2addr_b64 v[4:7], v12 offset0:62 offset1:63
	s_waitcnt lgkmcnt(0)
	v_add_f64 v[10:11], v[2:3], -v[4:5]
	v_add_f64 v[8:9], v[0:1], -v[6:7]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[10:11], v[10:11], s[4:5]
	v_fma_f64 v[20:21], v[8:9], s[2:3], -v[10:11]
	ds_load_2addr_b64 v[8:11], v12 offset0:67 offset1:68
	ds_load_2addr_b64 v[12:15], v12 offset0:60 offset1:61
	scratch_load_b32 v24, off, off offset:48 ; 4-byte Folded Reload
	s_waitcnt vmcnt(0)
	ds_load_2addr_b64 v[36:39], v24 offset0:80 offset1:96
	ds_load_2addr_b64 v[40:43], v24 offset0:32 offset1:48
	s_waitcnt lgkmcnt(2)
	v_add_f64 v[22:23], v[8:9], -v[14:15]
	ds_load_2addr_b64 v[64:67], v24 offset0:112 offset1:128
	s_clause 0x1
	scratch_store_b64 off, v[104:105], off offset:4
	scratch_store_b64 off, v[71:72], off offset:12
	ds_load_2addr_b64 v[68:71], v24 offset1:16
	v_dual_mov_b32 v105, v74 :: v_dual_mov_b32 v104, v73
	v_fma_f64 v[20:21], v[22:23], s[6:7], v[20:21]
	v_add_f64 v[22:23], v[10:11], -v[12:13]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_fma_f64 v[20:21], v[22:23], s[12:13], v[20:21]
	s_waitcnt lgkmcnt(2)
	v_add_f64 v[22:23], v[38:39], -v[40:41]
	v_fma_f64 v[18:19], -v[20:21], s[46:47], v[18:19]
	v_add_f64 v[20:21], v[36:37], -v[42:43]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[22:23], v[22:23], s[4:5]
	v_fma_f64 v[20:21], v[20:21], s[2:3], -v[22:23]
	s_waitcnt lgkmcnt(0)
	v_add_f64 v[22:23], v[64:65], -v[70:71]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f64 v[20:21], v[22:23], s[6:7], v[20:21]
	v_add_f64 v[22:23], v[66:67], -v[68:69]
	v_fma_f64 v[20:21], v[22:23], s[12:13], v[20:21]
	v_add_f64 v[22:23], v[78:79], -v[96:97]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_fma_f64 v[18:19], -v[20:21], s[48:49], v[18:19]
	v_add_f64 v[20:21], v[82:83], -v[100:101]
	v_mul_f64 v[22:23], v[22:23], s[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f64 v[20:21], v[20:21], s[2:3], -v[22:23]
	v_add_f64 v[22:23], v[80:81], -v[94:95]
	v_fma_f64 v[20:21], v[22:23], s[6:7], v[20:21]
	v_add_f64 v[22:23], v[150:151], -v[86:87]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[20:21], v[22:23], s[12:13], v[20:21]
	v_fma_f64 v[18:19], -v[20:21], s[50:51], v[18:19]
	global_store_b64 v[16:17], v[18:19], off
	ds_load_2addr_b64 v[16:19], v130 offset0:65 offset1:66
	ds_load_2addr_b64 v[44:47], v130 offset0:62 offset1:63
	ds_load_2addr_b64 v[48:51], v189 offset0:65 offset1:66
	ds_load_2addr_b64 v[52:55], v189 offset0:62 offset1:63
	ds_load_2addr_b64 v[20:23], v130 offset0:67 offset1:68
	ds_load_2addr_b64 v[28:31], v130 offset0:60 offset1:61
	ds_load_2addr_b64 v[24:27], v189 offset0:67 offset1:68
	ds_load_2addr_b64 v[32:35], v189 offset0:60 offset1:61
	s_waitcnt lgkmcnt(6)
	v_mul_f64 v[4:5], v[4:5], v[44:45]
	v_mul_f64 v[6:7], v[6:7], v[46:47]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_fma_f64 v[2:3], v[2:3], v[18:19], -v[4:5]
	s_waitcnt lgkmcnt(4)
	v_add_f64 v[4:5], v[50:51], -v[52:53]
	v_fma_f64 v[0:1], v[0:1], v[16:17], -v[6:7]
	v_add_f64 v[6:7], v[48:49], -v[54:55]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_add_f64 v[2:3], v[2:3], v[4:5]
	s_waitcnt lgkmcnt(0)
	v_add_f64 v[4:5], v[24:25], -v[34:35]
	v_add_f64 v[0:1], v[0:1], v[6:7]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[2:3], v[2:3], s[4:5]
	v_fma_f64 v[0:1], v[0:1], s[2:3], -v[2:3]
	v_mul_f64 v[2:3], v[14:15], v[30:31]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[2:3], v[8:9], v[20:21], -v[2:3]
	v_add_f64 v[2:3], v[2:3], v[4:5]
	v_add_f64 v[4:5], v[26:27], -v[32:33]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f64 v[0:1], v[2:3], s[6:7], v[0:1]
	v_mul_f64 v[2:3], v[12:13], v[28:29]
	v_fma_f64 v[2:3], v[10:11], v[22:23], -v[2:3]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[2:3], v[2:3], v[4:5]
	v_fma_f64 v[0:1], v[2:3], s[12:13], v[0:1]
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f64 v[182:183], -v[0:1], s[46:47], v[58:59]
	ds_load_2addr_b64 v[0:3], v168 offset0:80 offset1:96
	ds_load_2addr_b64 v[8:11], v137 offset0:80 offset1:96
	ds_load_2addr_b64 v[4:7], v168 offset0:32 offset1:48
	ds_load_2addr_b64 v[12:15], v137 offset0:32 offset1:48
	s_waitcnt lgkmcnt(0)
	v_mul_f64 v[4:5], v[4:5], v[12:13]
	v_mul_f64 v[6:7], v[6:7], v[14:15]
	v_mul_f64 v[40:41], v[40:41], v[12:13]
	v_mul_f64 v[42:43], v[42:43], v[14:15]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fma_f64 v[2:3], v[2:3], v[10:11], -v[4:5]
	v_fma_f64 v[0:1], v[0:1], v[8:9], -v[6:7]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fma_f64 v[38:39], v[38:39], v[10:11], -v[40:41]
	v_fma_f64 v[36:37], v[36:37], v[8:9], -v[42:43]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[2:3], v[2:3], s[4:5]
	v_fma_f64 v[172:173], v[0:1], s[2:3], -v[2:3]
	ds_load_2addr_b64 v[58:61], v168 offset0:112 offset1:128
	ds_load_2addr_b64 v[0:3], v137 offset0:112 offset1:128
	ds_load_2addr_b64 v[72:75], v168 offset1:16
	ds_load_2addr_b64 v[4:7], v137 offset1:16
	s_waitcnt lgkmcnt(0)
	v_mul_f64 v[74:75], v[74:75], v[6:7]
	v_mul_f64 v[72:73], v[72:73], v[4:5]
	v_mul_f64 v[68:69], v[68:69], v[4:5]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[58:59], v[58:59], v[0:1], -v[74:75]
	v_fma_f64 v[60:61], v[60:61], v[2:3], -v[72:73]
	v_mul_f64 v[72:73], v[118:119], v[124:125]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fma_f64 v[66:67], v[66:67], v[2:3], -v[68:69]
	v_fma_f64 v[58:59], v[58:59], s[6:7], v[172:173]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[72:73], v[138:139], v[92:93], -v[72:73]
	v_fma_f64 v[58:59], v[60:61], s[12:13], v[58:59]
	v_mul_f64 v[60:61], v[122:123], v[126:127]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_f64 v[72:73], v[72:73], s[4:5]
	v_fma_f64 v[58:59], -v[58:59], s[48:49], v[182:183]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[60:61], v[134:135], v[98:99], -v[60:61]
	v_fma_f64 v[60:61], v[60:61], s[2:3], -v[72:73]
	v_mul_f64 v[72:73], v[178:179], v[120:121]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[72:73], v[140:141], v[190:191], -v[72:73]
	v_fma_f64 v[60:61], v[72:73], s[6:7], v[60:61]
	v_mul_f64 v[72:73], v[174:175], v[170:171]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[72:73], v[88:89], v[166:167], -v[72:73]
	v_fma_f64 v[60:61], v[72:73], s[12:13], v[60:61]
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f64 v[58:59], -v[60:61], s[50:51], v[58:59]
	global_store_b64 v[62:63], v[58:59], off
	ds_load_2addr_b64 v[58:61], v186 offset0:65 offset1:66
	ds_load_2addr_b64 v[72:75], v186 offset0:62 offset1:63
	s_waitcnt lgkmcnt(0)
	v_mul_f64 v[62:63], v[46:47], v[74:75]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f64 v[58:59], v[16:17], v[58:59], -v[62:63]
	v_mul_f64 v[62:63], v[44:45], v[72:73]
	v_fma_f64 v[60:61], v[18:19], v[60:61], -v[62:63]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[60:61], v[60:61], s[4:5]
	v_fma_f64 v[62:63], v[58:59], s[2:3], -v[60:61]
	ds_load_2addr_b64 v[58:61], v186 offset0:67 offset1:68
	ds_load_2addr_b64 v[72:75], v186 offset0:60 offset1:61
	s_waitcnt lgkmcnt(0)
	v_mul_f64 v[74:75], v[30:31], v[74:75]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[58:59], v[20:21], v[58:59], -v[74:75]
	v_fma_f64 v[58:59], v[58:59], s[6:7], v[62:63]
	v_mul_f64 v[62:63], v[28:29], v[72:73]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[60:61], v[22:23], v[60:61], -v[62:63]
	v_fma_f64 v[58:59], v[60:61], s[12:13], v[58:59]
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f64 v[72:73], -v[58:59], s[46:47], v[56:57]
	ds_load_2addr_b64 v[56:59], v189 offset0:80 offset1:96
	ds_load_2addr_b64 v[60:63], v189 offset0:32 offset1:48
	s_waitcnt lgkmcnt(0)
	v_add_f64 v[40:41], v[58:59], -v[60:61]
	v_add_f64 v[42:43], v[56:57], -v[62:63]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[38:39], v[38:39], v[40:41]
	v_add_f64 v[36:37], v[36:37], v[42:43]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[38:39], v[38:39], s[4:5]
	v_fma_f64 v[74:75], v[36:37], s[2:3], -v[38:39]
	v_mul_f64 v[36:37], v[70:71], v[6:7]
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f64 v[64:65], v[64:65], v[0:1], -v[36:37]
	ds_load_2addr_b64 v[36:39], v189 offset0:112 offset1:128
	ds_load_2addr_b64 v[40:43], v189 offset1:16
	s_waitcnt lgkmcnt(0)
	v_add_f64 v[70:71], v[36:37], -v[42:43]
	v_add_f64 v[68:69], v[38:39], -v[40:41]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[64:65], v[64:65], v[70:71]
	v_add_f64 v[66:67], v[66:67], v[68:69]
	v_mul_f64 v[68:69], v[118:119], v[108:109]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[64:65], v[64:65], s[6:7], v[74:75]
	v_fma_f64 v[68:69], v[187:188], v[138:139], -v[68:69]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_fma_f64 v[64:65], v[66:67], s[12:13], v[64:65]
	v_mul_f64 v[66:67], v[84:85], v[122:123]
	v_mul_f64 v[68:69], v[68:69], s[4:5]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[64:65], -v[64:65], s[48:49], v[72:73]
	v_fma_f64 v[66:67], v[76:77], v[134:135], -v[66:67]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f64 v[66:67], v[66:67], s[2:3], -v[68:69]
	v_mul_f64 v[68:69], v[178:179], v[128:129]
	v_fma_f64 v[68:69], v[156:157], v[140:141], -v[68:69]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f64 v[66:67], v[68:69], s[6:7], v[66:67]
	v_mul_f64 v[68:69], v[174:175], v[142:143]
	v_fma_f64 v[68:69], v[116:117], v[88:89], -v[68:69]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[66:67], v[68:69], s[12:13], v[66:67]
	v_fma_f64 v[64:65], -v[66:67], s[50:51], v[64:65]
	global_store_b64 v[184:185], v[64:65], off
	ds_load_2addr_b64 v[64:67], v131 offset0:65 offset1:66
	ds_load_2addr_b64 v[68:71], v131 offset0:62 offset1:63
	s_waitcnt lgkmcnt(0)
	v_mul_f64 v[68:69], v[44:45], v[68:69]
	v_mul_f64 v[70:71], v[46:47], v[70:71]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[66:67], v[18:19], v[66:67], -v[68:69]
	v_fma_f64 v[64:65], v[16:17], v[64:65], -v[70:71]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[66:67], v[66:67], s[4:5]
	v_fma_f64 v[72:73], v[64:65], s[2:3], -v[66:67]
	ds_load_2addr_b64 v[64:67], v131 offset0:67 offset1:68
	ds_load_2addr_b64 v[68:71], v131 offset0:60 offset1:61
	s_waitcnt lgkmcnt(0)
	v_mul_f64 v[70:71], v[30:31], v[70:71]
	v_mul_f64 v[68:69], v[28:29], v[68:69]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[64:65], v[20:21], v[64:65], -v[70:71]
	v_fma_f64 v[66:67], v[22:23], v[66:67], -v[68:69]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[64:65], v[64:65], s[6:7], v[72:73]
	v_fma_f64 v[64:65], v[66:67], s[12:13], v[64:65]
	scratch_load_b64 v[66:67], off, off offset:36 ; 8-byte Folded Reload
	s_waitcnt vmcnt(0)
	v_fma_f64 v[72:73], -v[64:65], s[46:47], v[66:67]
	ds_load_2addr_b64 v[64:67], v131 offset0:80 offset1:96
	ds_load_2addr_b64 v[68:71], v131 offset0:32 offset1:48
	s_waitcnt lgkmcnt(0)
	v_mul_f64 v[68:69], v[12:13], v[68:69]
	v_mul_f64 v[70:71], v[14:15], v[70:71]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[66:67], v[10:11], v[66:67], -v[68:69]
	v_fma_f64 v[64:65], v[8:9], v[64:65], -v[70:71]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[66:67], v[66:67], s[4:5]
	v_fma_f64 v[74:75], v[64:65], s[2:3], -v[66:67]
	ds_load_2addr_b64 v[64:67], v131 offset0:112 offset1:128
	ds_load_2addr_b64 v[68:71], v131 offset1:16
	s_waitcnt lgkmcnt(0)
	v_mul_f64 v[70:71], v[6:7], v[70:71]
	v_mul_f64 v[68:69], v[4:5], v[68:69]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[64:65], v[0:1], v[64:65], -v[70:71]
	v_fma_f64 v[66:67], v[2:3], v[66:67], -v[68:69]
	v_add_f64 v[68:69], v[132:133], -v[148:149]
	v_add_f64 v[70:71], v[102:103], -v[146:147]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[64:65], v[64:65], s[6:7], v[74:75]
	v_fma_f64 v[64:65], v[66:67], s[12:13], v[64:65]
	v_mul_f64 v[66:67], v[100:101], v[122:123]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[64:65], -v[64:65], s[48:49], v[72:73]
	v_fma_f64 v[66:67], v[82:83], v[134:135], -v[66:67]
	v_dual_mov_b32 v73, v104 :: v_dual_mov_b32 v74, v105
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_f64 v[66:67], v[66:67], v[68:69]
	v_mul_f64 v[68:69], v[96:97], v[118:119]
	v_fma_f64 v[68:69], v[78:79], v[138:139], -v[68:69]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[68:69], v[68:69], v[70:71]
	v_add_f64 v[70:71], v[106:107], -v[144:145]
	v_mul_f64 v[68:69], v[68:69], s[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f64 v[66:67], v[66:67], s[2:3], -v[68:69]
	v_mul_f64 v[68:69], v[94:95], v[178:179]
	v_fma_f64 v[68:69], v[80:81], v[140:141], -v[68:69]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[68:69], v[68:69], v[70:71]
	v_add_f64 v[70:71], v[164:165], -v[90:91]
	v_fma_f64 v[66:67], v[68:69], s[6:7], v[66:67]
	v_mul_f64 v[68:69], v[86:87], v[174:175]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[68:69], v[150:151], v[88:89], -v[68:69]
	v_add_f64 v[68:69], v[68:69], v[70:71]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[66:67], v[68:69], s[12:13], v[66:67]
	v_fma_f64 v[64:65], -v[66:67], s[50:51], v[64:65]
	global_store_b64 v[180:181], v[64:65], off
	ds_load_2addr_b64 v[64:67], v136 offset0:65 offset1:66
	ds_load_2addr_b64 v[68:71], v136 offset0:62 offset1:63
	s_waitcnt lgkmcnt(0)
	v_mul_f64 v[70:71], v[46:47], v[70:71]
	v_mul_f64 v[46:47], v[46:47], v[54:55]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[64:65], v[16:17], v[64:65], -v[70:71]
	v_fma_f64 v[16:17], v[16:17], v[48:49], -v[46:47]
	v_mul_f64 v[46:47], v[44:45], v[68:69]
	v_mul_f64 v[44:45], v[44:45], v[52:53]
	s_clause 0x1
	scratch_load_b64 v[71:72], off, off offset:12
	scratch_load_b64 v[104:105], off, off offset:4
	v_add_f64 v[16:17], v[16:17], v[64:65]
	v_fma_f64 v[46:47], v[18:19], v[66:67], -v[46:47]
	v_fma_f64 v[18:19], v[18:19], v[50:51], -v[44:45]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[18:19], v[18:19], v[46:47]
	v_mul_f64 v[18:19], v[18:19], s[4:5]
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f64 v[48:49], v[16:17], s[2:3], -v[18:19]
	ds_load_2addr_b64 v[16:19], v136 offset0:67 offset1:68
	ds_load_2addr_b64 v[44:47], v136 offset0:60 offset1:61
	s_waitcnt lgkmcnt(0)
	v_mul_f64 v[46:47], v[30:31], v[46:47]
	v_mul_f64 v[30:31], v[30:31], v[34:35]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[16:17], v[20:21], v[16:17], -v[46:47]
	v_fma_f64 v[20:21], v[20:21], v[24:25], -v[30:31]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[16:17], v[20:21], v[16:17]
	v_mul_f64 v[20:21], v[28:29], v[44:45]
	v_fma_f64 v[16:17], v[16:17], s[6:7], v[48:49]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f64 v[18:19], v[22:23], v[18:19], -v[20:21]
	v_mul_f64 v[20:21], v[28:29], v[32:33]
	v_fma_f64 v[20:21], v[22:23], v[26:27], -v[20:21]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[18:19], v[20:21], v[18:19]
	v_fma_f64 v[16:17], v[18:19], s[12:13], v[16:17]
	scratch_load_b64 v[18:19], off, off offset:28 ; 8-byte Folded Reload
	s_waitcnt vmcnt(0)
	v_fma_f64 v[24:25], -v[16:17], s[46:47], v[18:19]
	ds_load_2addr_b64 v[16:19], v136 offset0:80 offset1:96
	ds_load_2addr_b64 v[20:23], v136 offset0:32 offset1:48
	s_waitcnt lgkmcnt(0)
	v_mul_f64 v[22:23], v[14:15], v[22:23]
	v_mul_f64 v[14:15], v[14:15], v[62:63]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[16:17], v[8:9], v[16:17], -v[22:23]
	v_fma_f64 v[8:9], v[8:9], v[56:57], -v[14:15]
	v_mul_f64 v[14:15], v[12:13], v[20:21]
	v_mul_f64 v[12:13], v[12:13], v[60:61]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[8:9], v[8:9], v[16:17]
	v_fma_f64 v[14:15], v[10:11], v[18:19], -v[14:15]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[10:11], v[10:11], v[58:59], -v[12:13]
	v_add_f64 v[10:11], v[10:11], v[14:15]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[10:11], v[10:11], s[4:5]
	v_fma_f64 v[16:17], v[8:9], s[2:3], -v[10:11]
	ds_load_2addr_b64 v[8:11], v136 offset0:112 offset1:128
	ds_load_2addr_b64 v[12:15], v136 offset1:16
	s_waitcnt lgkmcnt(0)
	v_mul_f64 v[14:15], v[6:7], v[14:15]
	v_mul_f64 v[6:7], v[6:7], v[42:43]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[8:9], v[0:1], v[8:9], -v[14:15]
	v_fma_f64 v[0:1], v[0:1], v[36:37], -v[6:7]
	v_mul_f64 v[6:7], v[4:5], v[12:13]
	v_mul_f64 v[4:5], v[4:5], v[40:41]
	scratch_load_b64 v[12:13], off, off offset:20 ; 8-byte Folded Reload
	v_add_f64 v[0:1], v[0:1], v[8:9]
	v_fma_f64 v[6:7], v[2:3], v[10:11], -v[6:7]
	v_fma_f64 v[2:3], v[2:3], v[38:39], -v[4:5]
	v_mul_f64 v[4:5], v[122:123], v[148:149]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[0:1], v[0:1], s[6:7], v[16:17]
	v_add_f64 v[2:3], v[2:3], v[6:7]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_fma_f64 v[4:5], v[134:135], v[132:133], -v[4:5]
	v_mul_f64 v[6:7], v[118:119], v[146:147]
	v_fma_f64 v[0:1], v[2:3], s[12:13], v[0:1]
	v_mul_f64 v[2:3], v[114:115], v[122:123]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[6:7], v[138:139], v[102:103], -v[6:7]
	v_fma_f64 v[0:1], -v[0:1], s[48:49], v[24:25]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[2:3], v[104:105], v[134:135], -v[2:3]
	v_add_f64 v[2:3], v[2:3], v[4:5]
	v_mul_f64 v[4:5], v[112:113], v[118:119]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[4:5], v[71:72], v[138:139], -v[4:5]
	v_add_f64 v[4:5], v[4:5], v[6:7]
	v_mul_f64 v[6:7], v[178:179], v[144:145]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f64 v[4:5], v[4:5], s[4:5]
	v_fma_f64 v[6:7], v[140:141], v[106:107], -v[6:7]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f64 v[2:3], v[2:3], s[2:3], -v[4:5]
	v_mul_f64 v[4:5], v[110:111], v[178:179]
	v_fma_f64 v[4:5], v[73:74], v[140:141], -v[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[4:5], v[4:5], v[6:7]
	v_mul_f64 v[6:7], v[174:175], v[90:91]
	v_fma_f64 v[2:3], v[4:5], s[6:7], v[2:3]
	v_mul_f64 v[4:5], v[154:155], v[174:175]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fma_f64 v[6:7], v[88:89], v[164:165], -v[6:7]
	s_waitcnt vmcnt(0)
	v_fma_f64 v[4:5], v[12:13], v[88:89], -v[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[4:5], v[4:5], v[6:7]
	v_fma_f64 v[2:3], v[4:5], s[12:13], v[2:3]
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f64 v[0:1], -v[2:3], s[50:51], v[0:1]
	global_store_b64 v[176:177], v[0:1], off
.LBB0_38:                               ;   in Loop: Header=BB0_12 Depth=1
	s_or_b32 exec_lo, exec_lo, s10
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_and_saveexec_b32 s2, s14
	s_cbranch_execz .LBB0_40
; %bb.39:                               ;   in Loop: Header=BB0_12 Depth=1
	ds_load_b64 v[0:1], v131 offset:512
	ds_load_b64 v[2:3], v186 offset:512
	ds_load_b64 v[4:5], v136 offset:512
	ds_store_b64 v186, v[76:77] offset:512
	ds_store_b64 v131, v[82:83] offset:512
	ds_store_b64 v136, v[104:105] offset:512
	ds_load_b64 v[6:7], v168 offset:512
	ds_load_b64 v[8:9], v189 offset:512
	ds_store_b64 v168, v[98:99] offset:512
	ds_store_b64 v189, v[132:133] offset:512
	v_dual_mov_b32 v76, v187 :: v_dual_mov_b32 v77, v188
	v_dual_mov_b32 v188, v157 :: v_dual_mov_b32 v187, v156
	v_dual_mov_b32 v157, v117 :: v_dual_mov_b32 v156, v116
	v_dual_mov_b32 v86, v94 :: v_dual_mov_b32 v87, v95
	v_dual_mov_b32 v94, v96 :: v_dual_mov_b32 v95, v97
	v_dual_mov_b32 v96, v100 :: v_dual_mov_b32 v97, v101
	v_dual_mov_b32 v83, v79 :: v_dual_mov_b32 v82, v78
	v_dual_mov_b32 v78, v80 :: v_dual_mov_b32 v79, v81
	v_dual_mov_b32 v80, v150 :: v_dual_mov_b32 v81, v151
	v_dual_mov_b32 v155, v111 :: v_dual_mov_b32 v154, v110
	s_waitcnt lgkmcnt(9)
	v_dual_mov_b32 v101, v1 :: v_dual_mov_b32 v100, v0
	v_dual_mov_b32 v110, v112 :: v_dual_mov_b32 v111, v113
	v_dual_mov_b32 v112, v114 :: v_dual_mov_b32 v113, v115
	s_waitcnt lgkmcnt(7)
	v_dual_mov_b32 v115, v5 :: v_dual_mov_b32 v114, v4
	v_dual_mov_b32 v105, v72 :: v_dual_mov_b32 v104, v71
	v_dual_mov_b32 v71, v73 :: v_dual_mov_b32 v72, v74
	v_dual_mov_b32 v74, v13 :: v_dual_mov_b32 v73, v12
	v_dual_mov_b32 v174, v178 :: v_dual_mov_b32 v175, v179
	v_dual_mov_b32 v179, v119 :: v_dual_mov_b32 v178, v118
	v_dual_mov_b32 v118, v122 :: v_dual_mov_b32 v119, v123
	v_dual_mov_b32 v122, v152 :: v_dual_mov_b32 v123, v153
	v_dual_mov_b32 v153, v135 :: v_dual_mov_b32 v152, v134
	v_dual_mov_b32 v134, v138 :: v_dual_mov_b32 v135, v139
	v_dual_mov_b32 v138, v140 :: v_dual_mov_b32 v139, v141
	v_dual_mov_b32 v141, v89 :: v_dual_mov_b32 v140, v88
	v_dual_mov_b32 v90, v144 :: v_dual_mov_b32 v91, v145
	v_dual_mov_b32 v144, v146 :: v_dual_mov_b32 v145, v147
	v_dual_mov_b32 v146, v148 :: v_dual_mov_b32 v147, v149
	s_waitcnt lgkmcnt(2)
	v_dual_mov_b32 v149, v9 :: v_dual_mov_b32 v148, v8
	v_dual_mov_b32 v133, v103 :: v_dual_mov_b32 v132, v102
	v_dual_mov_b32 v102, v106 :: v_dual_mov_b32 v103, v107
	v_dual_mov_b32 v106, v164 :: v_dual_mov_b32 v107, v165
	v_dual_mov_b32 v143, v129 :: v_dual_mov_b32 v142, v128
	v_dual_mov_b32 v99, v93 :: v_dual_mov_b32 v98, v92
	v_dual_mov_b32 v171, v121 :: v_dual_mov_b32 v170, v120
	v_dual_mov_b32 v120, v124 :: v_dual_mov_b32 v121, v125
	v_dual_mov_b32 v124, v126 :: v_dual_mov_b32 v125, v127
	v_dual_mov_b32 v127, v7 :: v_dual_mov_b32 v126, v6
	v_dual_mov_b32 v92, v190 :: v_dual_mov_b32 v93, v191
	v_dual_mov_b32 v191, v167 :: v_dual_mov_b32 v190, v166
	v_dual_mov_b32 v129, v109 :: v_dual_mov_b32 v128, v108
	v_dual_mov_b32 v109, v85 :: v_dual_mov_b32 v108, v84
	v_dual_mov_b32 v85, v3 :: v_dual_mov_b32 v84, v2
.LBB0_40:                               ;   in Loop: Header=BB0_12 Depth=1
	s_or_b32 exec_lo, exec_lo, s2
	v_add_nc_u32_e32 v158, s9, v158
	v_add_nc_u32_e32 v159, s9, v159
	v_add_nc_u32_e32 v160, s9, v160
	v_add_nc_u32_e32 v161, s9, v161
	v_add_nc_u32_e32 v162, s9, v162
	v_add_nc_u32_e32 v163, s9, v163
	s_add_i32 s8, s8, -1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_eq_u32 s8, 0
	s_cbranch_scc0 .LBB0_12
.LBB0_41:                               ; %._crit_edge
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii
		.amdhsa_group_segment_fixed_size 21504
		.amdhsa_private_segment_fixed_size 76
		.amdhsa_kernarg_size 400
		.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 1
		.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 192
		.amdhsa_next_free_sgpr 57
		.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	_Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii, .Lfunc_end0-_Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 7484
; NumSgprs: 59
; NumVgprs: 192
; ScratchSize: 76
; MemoryBound: 1
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 21504 bytes/workgroup (compile time only)
; SGPRBlocks: 7
; VGPRBlocks: 23
; NumSGPRsForWavesPerEU: 59
; NumVGPRsForWavesPerEU: 192
; Occupancy: 8
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 1
; 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
      - .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
      - .address_space:  global
        .offset:         48
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         56
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         64
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         72
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         80
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         88
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         96
        .size:           8
        .value_kind:     global_buffer
      - .offset:         104
        .size:           8
        .value_kind:     by_value
      - .offset:         112
        .size:           8
        .value_kind:     by_value
      - .offset:         120
        .size:           8
        .value_kind:     by_value
      - .offset:         128
        .size:           4
        .value_kind:     by_value
      - .offset:         132
        .size:           4
        .value_kind:     by_value
      - .offset:         136
        .size:           4
        .value_kind:     by_value
      - .offset:         144
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         148
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         152
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         156
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         158
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         160
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         162
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         164
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         166
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         184
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         192
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         200
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         208
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 21504
    .kernarg_segment_align: 8
    .kernarg_segment_size: 400
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii
    .private_segment_fixed_size: 76
    .sgpr_count:     59
    .sgpr_spill_count: 0
    .symbol:         _Z7hyptermPdS_S_S_S_S_S_S_S_S_S_S_S_dddiii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     192
    .vgpr_spill_count: 18
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
