__device__ int find_qmin_and_qmax(
        double dq0, 
        double dq1, 
        double dq2, 
        double *qmin, 
        double *qmax)
{
    // Considering the centroid of an FV triangle and the vertices of its 
    // auxiliary triangle, find 
    // qmin=min(q)-qc and qmax=max(q)-qc, 
    // where min(q) and max(q) are respectively min and max over the
    // four values (at the centroid of the FV triangle and the auxiliary 
    // triangle vertices),
    // and qc is the centroid
    // dq0=q(vertex0)-q(centroid of FV triangle)
    // dq1=q(vertex1)-q(vertex0)
    // dq2=q(vertex2)-q(vertex0)

    // This is a simple implementation 
    *qmax = max(max(dq0, max(dq0+dq1, dq0+dq2)), 0.0) ;
    *qmin = min(min(dq0, min(dq0+dq1, dq0+dq2)), 0.0) ;

    return 0;
}



__device__ int limit_gradient(
        double *dqv, 
        double qmin, 
        double qmax, 
        double beta_w)
{
    // Given provisional jumps dqv from the FV triangle centroid to its 
    // vertices and jumps qmin (qmax) between the centroid of the FV 
    // triangle and the minimum (maximum) of the values at the centroid of 
    // the FV triangle and the auxiliary triangle vertices,
    // calculate a multiplicative factor phi by which the provisional 
    // vertex jumps are to be limited

    int i;
    double r=1000.0, r0=1.0, phi=1.0;
    double TINY = 1.0e-100; // to avoid machine accuracy problems.
    // FIXME: Perhaps use the epsilon used elsewhere.

    // Any provisional jump with magnitude < TINY does not contribute to 
    // the limiting process.

    for (i=0;i<3;i++){
        if (dqv[i]<-TINY)
            r0=qmin/dqv[i];

        if (dqv[i]>TINY)
            r0=qmax/dqv[i];

        r=min(r0,r);
    }

    phi=min(r*beta_w,1.0);
    //phi=1.;
    dqv[0]=dqv[0]*phi;
    dqv[1]=dqv[1]*phi;
    dqv[2]=dqv[2]*phi;

    return 0;
}



// Computational routine
__global__ void  _extrapolate_second_order_edge_sw(
        //int number_of_elements,
        int N,
        int optimise_dry_cells, 
        int extrapolate_velocity_second_order,
        double epsilon,
        double minimum_allowed_height,
        double beta_w,
        double beta_w_dry,
        double beta_uh,
        double beta_uh_dry,
        double beta_vh,
        double beta_vh_dry,

        long* surrogate_neighbours,
        long* number_of_boundaries,
        double* centroid_coordinates,

        double* stage_centroid_values,
        double* elevation_centroid_values,
        double* xmom_centroid_values,
        double* ymom_centroid_values,

        double* edge_coordinates,

        double* stage_edge_values,
        double* elevation_edge_values,
        double* xmom_edge_values,
        double* ymom_edge_values,

        double* stage_vertex_values,
        double* elevation_vertex_values,
        double* xmom_vertex_values,
        double* ymom_vertex_values,

        double * stage_centroid_store,
        double * xmom_centroid_store,
        double * ymom_centroid_store,
        double * min_elevation_edgevalue,
        double * max_elevation_edgevalue,
        int * count_wet_neighbours
        ) 
{
    const int k = 
        threadIdx.x+threadIdx.y*blockDim.x+
        (blockIdx.x+blockIdx.y*gridDim.x)*blockDim.x*blockDim.y;

    if ( k >= N)
        return;
    // Local variables
    // Gradient vector used to calculate edge values from centroids
    double a, b; 
    int k0, k1, k2, k3, k6, coord_index, i, ii, ktmp;
    // Vertices of the auxiliary triangle
    double x, y, x0, y0, x1, y1, x2, y2, xv0, yv0, xv1, yv1, xv2, yv2; 
    double dx1, dx2, dy1, dy2;
    double dxv0, dxv1, dxv2, dyv0, dyv1, dyv2, dq0, dq1, dq2;
    double area2, inv_area2;
    double dqv[3], qmin, qmax, hmin, hmax, bedmax, stagemin;
    double hc, h0, h1, h2, beta_tmp, hfactor, xtmp, ytmp;
    double dk, dv0, dv1, dv2, de[3];
    double demin, dcmax, r0scale, vect_norm, l1, l2;

    //double *xmom_centroid_store, *ymom_centroid_store; 
    //double *stage_centroid_store, *min_elevation_edgevalue;
    //double *max_elevation_edgevalue;
    //int *count_wet_neighbours;

    // Use malloc to avoid putting these variables on the stack, 
    // which can cause
    // segfaults in large model runs
    //xmom_centroid_store = malloc(number_of_elements*sizeof(double));
    //ymom_centroid_store = malloc(number_of_elements*sizeof(double));
    //stage_centroid_store = malloc(number_of_elements*sizeof(double));
    //min_elevation_edgevalue = malloc(number_of_elements*sizeof(double));
    //max_elevation_edgevalue = malloc(number_of_elements*sizeof(double));
    //count_wet_neighbours = malloc(number_of_elements*sizeof(int));

    if(extrapolate_velocity_second_order==1){
        // Replace momentum centroid with velocity centroid 2 allow velocity
        // extrapolation This will be changed back at the end of the routine
        //for (k=0; k<number_of_elements; k++){

        dk = max(stage_centroid_values[k]-elevation_centroid_values[k],
                minimum_allowed_height);
        xmom_centroid_store[k] = xmom_centroid_values[k];
        xmom_centroid_values[k] = xmom_centroid_values[k]/dk;

        ymom_centroid_store[k] = ymom_centroid_values[k];
        ymom_centroid_values[k] = ymom_centroid_values[k]/dk;

        min_elevation_edgevalue[k] = min(elevation_edge_values[3*k], 
                min(elevation_edge_values[3*k+1],
                    elevation_edge_values[3*k+2]));
        max_elevation_edgevalue[k] = max(elevation_edge_values[3*k], 
                max(elevation_edge_values[3*k+1],
                    elevation_edge_values[3*k+2]));
        //} // for k=0 to number_of_elements-1
    }

    // Count how many 'fully submerged' neighbours the cell has
    //for(k=0; k<number_of_elements;k++){ 
    count_wet_neighbours[k]=0;
    for (i=0; i<3; i++){
        ktmp = surrogate_neighbours[3*k+i];              
        if(stage_centroid_values[ktmp] > max_elevation_edgevalue[ktmp]){
            count_wet_neighbours[k]+=1;
        }
    }
    //} // for k=0 to number_of_elements-1

    // Begin extrapolation routine
    //for (k = 0; k < number_of_elements; k++) {
    do{
    k3=k*3;
    k6=k*6;

    if (number_of_boundaries[k]==3)
        //if (0==0)
    {
        // No neighbours, set gradient on the triangle to zero

        stage_edge_values[k3]   = stage_centroid_values[k];
        stage_edge_values[k3+1] = stage_centroid_values[k];
        stage_edge_values[k3+2] = stage_centroid_values[k];
        xmom_edge_values[k3]    = xmom_centroid_values[k];
        xmom_edge_values[k3+1]  = xmom_centroid_values[k];
        xmom_edge_values[k3+2]  = xmom_centroid_values[k];
        ymom_edge_values[k3]    = ymom_centroid_values[k];
        ymom_edge_values[k3+1]  = ymom_centroid_values[k];
        ymom_edge_values[k3+2]  = ymom_centroid_values[k];

        continue;
    }
    else 
    {
        // Triangle k has one or more neighbours. 
        // Get centroid and edge coordinates of the triangle

        // Get the edge coordinates
        xv0 = edge_coordinates[k6];   
        yv0 = edge_coordinates[k6+1];
        xv1 = edge_coordinates[k6+2]; 
        yv1 = edge_coordinates[k6+3];
        xv2 = edge_coordinates[k6+4]; 
        yv2 = edge_coordinates[k6+5];

        // Get the centroid coordinates
        coord_index = 2*k;
        x = centroid_coordinates[coord_index];
        y = centroid_coordinates[coord_index+1];

        // Store x- and y- differentials for the edges of 
        // triangle k relative to the centroid
        dxv0 = xv0 - x; 
        dxv1 = xv1 - x; 
        dxv2 = xv2 - x;
        dyv0 = yv0 - y; 
        dyv1 = yv1 - y; 
        dyv2 = yv2 - y;
        // Compute the minimum distance from the centroid to an edge
        //demin=min(dxv0*dxv0 +dyv0*dyv0, 
        //      min(dxv1*dxv1+dyv1*dyv1, dxv2*dxv2+dyv2*dyv2));
        //demin=sqrt(demin);
    }



    if (number_of_boundaries[k]<=1)
    {
        //==============================================
        // Number of boundaries <= 1
        //==============================================    


        // If no boundaries, auxiliary triangle is formed 
        // from the centroids of the three neighbours
        // If one boundary, auxiliary triangle is formed 
        // from this centroid and its two neighbours

        k0 = surrogate_neighbours[k3];
        k1 = surrogate_neighbours[k3 + 1];
        k2 = surrogate_neighbours[k3 + 2];

        // Take note if the max neighbour bed elevation 
        // is greater than the min
        // neighbour stage -- suggests a 'steep' bed relative to the flow
        bedmax = max(elevation_centroid_values[k], 
                    max(elevation_centroid_values[k0],
                        max(elevation_centroid_values[k1], 
                            elevation_centroid_values[k2])));
        //bedmax = elevation_centroid_values[k];
        stagemin = min(max(stage_centroid_values[k],
                            elevation_centroid_values[k]), 
                min(max(stage_centroid_values[k0], 
                        elevation_centroid_values[k0]),
                    min(max(stage_centroid_values[k1], 
                        elevation_centroid_values[k1]),
                        max(stage_centroid_values[k2], 
                            elevation_centroid_values[k2]))));
        if(stagemin < bedmax){
            // This will cause first order extrapolation
            k2 = k;
            k0 = k;
            k1 = k;
        } 

        // Get the auxiliary triangle's vertex coordinates 
        // (really the centroids of neighbouring triangles)
        coord_index = 2*k0;
        x0 = centroid_coordinates[coord_index];
        y0 = centroid_coordinates[coord_index+1];

        coord_index = 2*k1;
        x1 = centroid_coordinates[coord_index];
        y1 = centroid_coordinates[coord_index+1];

        coord_index = 2*k2;
        x2 = centroid_coordinates[coord_index];
        y2 = centroid_coordinates[coord_index+1];

        // Store x- and y- differentials for the vertices 
        // of the auxiliary triangle
        dx1 = x1 - x0; 
        dx2 = x2 - x0;
        dy1 = y1 - y0; 
        dy2 = y2 - y0;

        // Calculate 2*area of the auxiliary triangle
        // The triangle is guaranteed to be counter-clockwise      
        area2 = dy2*dx1 - dy1*dx2; 

        // Treat triangles with zero or 1 wet neighbours. 
        if ((area2 <= 0)) //|(count_wet_neighbours[k]==0))
        {
            //printf("Error negative triangle area \n");
            //report_python_error(AT, "Negative triangle area");
            //return -1;
            stage_edge_values[k3]   = stage_centroid_values[k];
            stage_edge_values[k3+1] = stage_centroid_values[k];
            stage_edge_values[k3+2] = stage_centroid_values[k];
            // First order momentum / velocity extrapolation
            xmom_edge_values[k3]    = xmom_centroid_values[k];
            xmom_edge_values[k3+1]  = xmom_centroid_values[k];
            xmom_edge_values[k3+2]  = xmom_centroid_values[k];
            ymom_edge_values[k3]    = ymom_centroid_values[k];
            ymom_edge_values[k3+1]  = ymom_centroid_values[k];
            ymom_edge_values[k3+2]  = ymom_centroid_values[k];

            continue;
        }  

        // Calculate heights of neighbouring cells
        hc = stage_centroid_values[k]  - elevation_centroid_values[k];
        h0 = stage_centroid_values[k0] - elevation_centroid_values[k0];
        h1 = stage_centroid_values[k1] - elevation_centroid_values[k1];
        h2 = stage_centroid_values[k2] - elevation_centroid_values[k2];
        hmin = min(min(h0, min(h1, h2)), hc);

        hfactor = 0.0;
        //if (hmin > 0.001) 
        if (hmin > 0.) 
            //if (hc>0.0)
        {
            hfactor = 1.0 ;//hmin/(hmin + 0.004);
            //hfactor=hmin/(hmin + 0.004);
        }


        //-----------------------------------
        // stage
        //-----------------------------------      

        // Calculate the difference between vertex 0 of the auxiliary 
        // triangle and the centroid of triangle k
        dq0 = stage_centroid_values[k0] - stage_centroid_values[k];

        // Calculate differentials between the vertices 
        // of the auxiliary triangle (centroids of neighbouring triangles)
        dq1 = stage_centroid_values[k1] - stage_centroid_values[k0];
        dq2 = stage_centroid_values[k2] - stage_centroid_values[k0];

        inv_area2 = 1.0/area2;
        // Calculate the gradient of stage on the auxiliary triangle
        a = dy2*dq1 - dy1*dq2;
        a *= inv_area2;
        b = dx1*dq2 - dx2*dq1;
        b *= inv_area2;

        // Calculate provisional jumps in stage from the centroid 
        // of triangle k to its vertices, to be limited
        dqv[0] = a*dxv0 + b*dyv0;
        dqv[1] = a*dxv1 + b*dyv1;
        dqv[2] = a*dxv2 + b*dyv2;

        // Now we want to find min and max of the centroid and the 
        // vertices of the auxiliary triangle and compute jumps 
        // from the centroid to the min and max
        find_qmin_and_qmax(dq0, dq1, dq2, &qmin, &qmax);

        beta_tmp = beta_w_dry + (beta_w - beta_w_dry) * hfactor;


        // Limit the gradient
        limit_gradient(dqv, qmin, qmax, beta_tmp);
        stage_edge_values[k3+0] = stage_centroid_values[k] + dqv[0];
        stage_edge_values[k3+1] = stage_centroid_values[k] + dqv[1];
        stage_edge_values[k3+2] = stage_centroid_values[k] + dqv[2];


        //-----------------------------------
        // xmomentum
        //-----------------------------------            

        // Calculate the difference between vertex 0 of the auxiliary 
        // triangle and the centroid of triangle k      
        dq0 = xmom_centroid_values[k0] - xmom_centroid_values[k];

        // Calculate differentials between the vertices 
        // of the auxiliary triangle
        dq1 = xmom_centroid_values[k1] - xmom_centroid_values[k0];
        dq2 = xmom_centroid_values[k2] - xmom_centroid_values[k0];

        // Calculate the gradient of xmom on the auxiliary triangle
        a = dy2*dq1 - dy1*dq2;
        a *= inv_area2;
        b = dx1*dq2 - dx2*dq1;
        b *= inv_area2;

        // Calculate provisional jumps in stage from the centroid 
        // of triangle k to its vertices, to be limited      
        dqv[0] = a*dxv0+b*dyv0;
        dqv[1] = a*dxv1+b*dyv1;
        dqv[2] = a*dxv2+b*dyv2;

        // Now we want to find min and max of the centroid and the 
        // vertices of the auxiliary triangle and compute jumps 
        // from the centroid to the min and max
        //
        find_qmin_and_qmax(dq0, dq1, dq2, &qmin, &qmax);

        beta_tmp = beta_uh_dry + (beta_uh - beta_uh_dry) * hfactor;

        limit_gradient(dqv, qmin, qmax, beta_tmp);


        for (i=0; i < 3; i++)
        {
            xmom_edge_values[k3+i] = xmom_centroid_values[k] + dqv[i];
        }

        //-----------------------------------
        // ymomentum
        //-----------------------------------                  

        // Calculate the difference between vertex 0 of the auxiliary 
        // triangle and the centroid of triangle k
        dq0 = ymom_centroid_values[k0] - ymom_centroid_values[k];

        // Calculate differentials between the vertices 
        // of the auxiliary triangle
        dq1 = ymom_centroid_values[k1] - ymom_centroid_values[k0];
        dq2 = ymom_centroid_values[k2] - ymom_centroid_values[k0];

        // Calculate the gradient of xmom on the auxiliary triangle
        a = dy2*dq1 - dy1*dq2;
        a *= inv_area2;
        b = dx1*dq2 - dx2*dq1;
        b *= inv_area2;

        // Calculate provisional jumps in stage from the centroid 
        // of triangle k to its vertices, to be limited
        dqv[0] = a*dxv0 + b*dyv0;
        dqv[1] = a*dxv1 + b*dyv1;
        dqv[2] = a*dxv2 + b*dyv2;

        // Now we want to find min and max of the centroid and the 
        // vertices of the auxiliary triangle and compute jumps 
        // from the centroid to the min and max
        //
        find_qmin_and_qmax(dq0, dq1, dq2, &qmin, &qmax);

        beta_tmp = beta_vh_dry + (beta_vh - beta_vh_dry) * hfactor;   

        limit_gradient(dqv, qmin, qmax, beta_tmp);

        for (i=0;i<3;i++)
        {
            ymom_edge_values[k3 + i] = ymom_centroid_values[k] + dqv[i];
        }

    } // End number_of_boundaries <=1 
    else
    {

        //==============================================
        // Number of boundaries == 2
        //==============================================        

        // One internal neighbour and gradient 
        // is in direction of the neighbour's centroid

        // Find the only internal neighbour (k1?)
        for (k2 = k3; k2 < k3 + 3; k2++)
        {
            // Find internal neighbour of triangle k      
            // k2 indexes the edges of triangle k   

            if (surrogate_neighbours[k2] != k)
            {
                break;
            }
        }

        if ((k2 == k3 + 3)) 
        {
            // If we didn't find an internal neighbour
            //report_python_error(AT, "Internal neighbour not found");
            //return -1;
            return;
        }

        k1 = surrogate_neighbours[k2];

        // The coordinates of the triangle are already (x,y). 
        // Get centroid of the neighbour (x1,y1)
        coord_index = 2*k1;
        x1 = centroid_coordinates[coord_index];
        y1 = centroid_coordinates[coord_index + 1];

        // Compute x- and y- distances between the centroid of 
        // triangle k and that of its neighbour
        dx1 = x1 - x; 
        dy1 = y1 - y;

        // Set area2 as the square of the distance
        area2 = dx1*dx1 + dy1*dy1;

        // Set dx2=(x1-x0)/((x1-x0)^2+(y1-y0)^2) 
        // and dy2=(y1-y0)/((x1-x0)^2+(y1-y0)^2) which
        // respectively correspond to the x- and y- gradients 
        // of the conserved quantities
        dx2 = 1.0/area2;
        dy2 = dx2*dy1;
        dx2 *= dx1;


        //-----------------------------------
        // stage
        //-----------------------------------            

        // Compute differentials
        dq1 = stage_centroid_values[k1] - stage_centroid_values[k];

        // Calculate the gradient between the centroid of triangle k 
        // and that of its neighbour
        a = dq1*dx2;
        b = dq1*dy2;

        // Calculate provisional edge jumps, to be limited
        dqv[0] = a*dxv0 + b*dyv0;
        dqv[1] = a*dxv1 + b*dyv1;
        dqv[2] = a*dxv2 + b*dyv2;

        // Now limit the jumps
        if (dq1>=0.0)
        {
            qmin=0.0;
            qmax=dq1;
        }
        else
        {
            qmin = dq1;
            qmax = 0.0;
        }

        // Limit the gradient
        limit_gradient(dqv, qmin, qmax, beta_w);

        //for (i=0; i < 3; i++)
        //{
        stage_edge_values[k3] = stage_centroid_values[k] + dqv[0];
        stage_edge_values[k3 + 1] = stage_centroid_values[k] + dqv[1];
        stage_edge_values[k3 + 2] = stage_centroid_values[k] + dqv[2];
        //}

        //-----------------------------------
        // xmomentum
        //-----------------------------------                        

        // Compute differentials
        dq1 = xmom_centroid_values[k1] - xmom_centroid_values[k];

        // Calculate the gradient between the centroid of triangle k 
        // and that of its neighbour
        a = dq1*dx2;
        b = dq1*dy2;

        // Calculate provisional edge jumps, to be limited
        dqv[0] = a*dxv0+b*dyv0;
        dqv[1] = a*dxv1+b*dyv1;
        dqv[2] = a*dxv2+b*dyv2;

        // Now limit the jumps
        if (dq1 >= 0.0)
        {
            qmin = 0.0;
            qmax = dq1;
        }
        else
        {
            qmin = dq1;
            qmax = 0.0;
        }

        // Limit the gradient      
        limit_gradient(dqv, qmin, qmax, beta_w);

        //for (i=0;i<3;i++)
        //xmom_edge_values[k3] = xmom_centroid_values[k] + dqv[0];
        //xmom_edge_values[k3 + 1] = xmom_centroid_values[k] + dqv[1];
        //xmom_edge_values[k3 + 2] = xmom_centroid_values[k] + dqv[2];

        for (i = 0; i < 3;i++)
        {
            xmom_edge_values[k3 + i] = xmom_centroid_values[k] + dqv[i];
        }

        //-----------------------------------
        // ymomentum
        //-----------------------------------                        

        // Compute differentials
        dq1 = ymom_centroid_values[k1] - ymom_centroid_values[k];

        // Calculate the gradient between the centroid of triangle k 
        // and that of its neighbours      a = dq1*dx2;
        b = dq1*dy2;

        // Calculate provisional edge jumps, to be limited
        dqv[0] = a*dxv0 + b*dyv0;
        dqv[1] = a*dxv1 + b*dyv1;
        dqv[2] = a*dxv2 + b*dyv2;

        // Now limit the jumps
        if (dq1>=0.0)
        {
            qmin = 0.0;
            qmax = dq1;
        } 
        else
        {
            qmin = dq1;
            qmax = 0.0;
        }

        // Limit the gradient
        limit_gradient(dqv, qmin, qmax, beta_w);

        for (i=0;i<3;i++)
        {
            ymom_edge_values[k3 + i] = ymom_centroid_values[k] + dqv[i];
        }
    } // else [number_of_boundaries==2]
    } while(k < 0);
    //} // for k=0 to number_of_elements-1

    // Compute vertex values of quantities
    //for (k=0; k<number_of_elements; k++){
    k3=3*k;

    // Compute stage vertex values 
    stage_vertex_values[k3] = stage_edge_values[k3+1] + stage_edge_values[k3+2] -stage_edge_values[k3] ; 
    stage_vertex_values[k3+1] =  stage_edge_values[k3] + stage_edge_values[k3+2]-stage_edge_values[k3+1]; 
    stage_vertex_values[k3+2] =  stage_edge_values[k3] + stage_edge_values[k3+1]-stage_edge_values[k3+2]; 

    // Compute xmom vertex values 
    xmom_vertex_values[k3] = xmom_edge_values[k3+1] + xmom_edge_values[k3+2] -xmom_edge_values[k3] ; 
    xmom_vertex_values[k3+1] =  xmom_edge_values[k3] + xmom_edge_values[k3+2]-xmom_edge_values[k3+1]; 
    xmom_vertex_values[k3+2] =  xmom_edge_values[k3] + xmom_edge_values[k3+1]-xmom_edge_values[k3+2]; 

    // Compute ymom vertex values 
    ymom_vertex_values[k3] = ymom_edge_values[k3+1] + ymom_edge_values[k3+2] -ymom_edge_values[k3] ; 
    ymom_vertex_values[k3+1] =  ymom_edge_values[k3] + ymom_edge_values[k3+2]-ymom_edge_values[k3+1]; 
    ymom_vertex_values[k3+2] =  ymom_edge_values[k3] + ymom_edge_values[k3+1]-ymom_edge_values[k3+2]; 

    // If needed, convert from velocity to momenta
    if(extrapolate_velocity_second_order==1){
        //Convert velocity back to momenta at centroids
        xmom_centroid_values[k] = xmom_centroid_store[k];
        ymom_centroid_values[k] = ymom_centroid_store[k];

        // Re-compute momenta at edges
        for (i=0; i<3; i++){
            de[i] = max(stage_edge_values[k3+i]-elevation_edge_values[k3+i],0.0);
            xmom_edge_values[k3+i]=xmom_edge_values[k3+i]*de[i];
            ymom_edge_values[k3+i]=ymom_edge_values[k3+i]*de[i];
        }

        // Re-compute momenta at vertices
        for (i=0; i<3; i++){
            de[i] = max(stage_vertex_values[k3+i]-elevation_vertex_values[k3+i],0.0);
            xmom_vertex_values[k3+i]=xmom_vertex_values[k3+i]*de[i];
            ymom_vertex_values[k3+i]=ymom_vertex_values[k3+i]*de[i];
        }

    }
    //} // for k=0 to number_of_elements-1

}  

	.file	"tmpxft_002f4520_00000000-6_extrapolate_second_order_edge_swb2.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2012:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2012:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z18find_qmin_and_qmaxdddPdS_
	.type	_Z18find_qmin_and_qmaxdddPdS_, @function
_Z18find_qmin_and_qmaxdddPdS_:
.LFB2008:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2008:
	.size	_Z18find_qmin_and_qmaxdddPdS_, .-_Z18find_qmin_and_qmaxdddPdS_
	.globl	_Z14limit_gradientPdddd
	.type	_Z14limit_gradientPdddd, @function
_Z14limit_gradientPdddd:
.LFB2009:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2009:
	.size	_Z14limit_gradientPdddd, .-_Z14limit_gradientPdddd
	.globl	_Z124__device_stub__Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_PiiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi
	.type	_Z124__device_stub__Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_PiiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi, @function
_Z124__device_stub__Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_PiiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi:
.LFB2034:
	.cfi_startproc
	endbr64
	subq	$600, %rsp
	.cfi_def_cfa_offset 608
	movq	608(%rsp), %rax
	movl	%edi, 252(%rsp)
	leaq	272(%rsp), %rdi
	movl	%esi, 248(%rsp)
	leaq	284(%rsp), %rsi
	movq	%rax, 144(%rsp)
	movq	616(%rsp), %rax
	movl	%edx, 244(%rsp)
	leaq	256(%rsp), %rdx
	movq	%rax, 136(%rsp)
	movq	624(%rsp), %rax
	movq	%rcx, 168(%rsp)
	leaq	264(%rsp), %rcx
	movq	%rax, 128(%rsp)
	movq	632(%rsp), %rax
	movq	%r8, 160(%rsp)
	movq	%rax, 120(%rsp)
	movq	640(%rsp), %rax
	movq	%r9, 152(%rsp)
	movq	%rax, 112(%rsp)
	movq	648(%rsp), %rax
	movsd	%xmm0, 232(%rsp)
	movq	%rax, 104(%rsp)
	movq	656(%rsp), %rax
	movsd	%xmm1, 224(%rsp)
	movq	%rax, 96(%rsp)
	movq	664(%rsp), %rax
	movsd	%xmm2, 216(%rsp)
	movq	%rax, 88(%rsp)
	movsd	%xmm3, 208(%rsp)
	movsd	%xmm4, 200(%rsp)
	movsd	%xmm5, 192(%rsp)
	movsd	%xmm6, 184(%rsp)
	movsd	%xmm7, 176(%rsp)
	movq	672(%rsp), %rax
	movq	%rax, 80(%rsp)
	movq	680(%rsp), %rax
	movq	%rax, 72(%rsp)
	movq	688(%rsp), %rax
	movq	%rax, 64(%rsp)
	movq	696(%rsp), %rax
	movq	%rax, 56(%rsp)
	movq	704(%rsp), %rax
	movq	%rax, 48(%rsp)
	movq	712(%rsp), %rax
	movq	%rax, 40(%rsp)
	movq	720(%rsp), %rax
	movq	%rax, 32(%rsp)
	movq	728(%rsp), %rax
	movq	%rax, 24(%rsp)
	movq	736(%rsp), %rax
	movq	%rax, 16(%rsp)
	movq	744(%rsp), %rax
	movq	%rax, 8(%rsp)
	movq	752(%rsp), %rax
	movq	%rax, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 584(%rsp)
	xorl	%eax, %eax
	leaq	252(%rsp), %rax
	movq	%rax, 320(%rsp)
	leaq	248(%rsp), %rax
	movq	%rax, 328(%rsp)
	leaq	244(%rsp), %rax
	movq	%rax, 336(%rsp)
	leaq	232(%rsp), %rax
	movq	%rax, 344(%rsp)
	leaq	224(%rsp), %rax
	movq	%rax, 352(%rsp)
	leaq	216(%rsp), %rax
	movq	%rax, 360(%rsp)
	leaq	208(%rsp), %rax
	movq	%rax, 368(%rsp)
	leaq	200(%rsp), %rax
	movq	%rax, 376(%rsp)
	leaq	192(%rsp), %rax
	movq	%rax, 384(%rsp)
	leaq	184(%rsp), %rax
	movq	%rax, 392(%rsp)
	leaq	176(%rsp), %rax
	movq	%rax, 400(%rsp)
	leaq	168(%rsp), %rax
	movq	%rax, 408(%rsp)
	leaq	160(%rsp), %rax
	movq	%rax, 416(%rsp)
	leaq	152(%rsp), %rax
	movq	%rax, 424(%rsp)
	leaq	144(%rsp), %rax
	movq	%rax, 432(%rsp)
	leaq	136(%rsp), %rax
	movq	%rax, 440(%rsp)
	leaq	128(%rsp), %rax
	movq	%rax, 448(%rsp)
	leaq	120(%rsp), %rax
	movq	%rax, 456(%rsp)
	leaq	112(%rsp), %rax
	movq	%rax, 464(%rsp)
	leaq	104(%rsp), %rax
	movq	%rax, 472(%rsp)
	leaq	96(%rsp), %rax
	movq	%rax, 480(%rsp)
	leaq	88(%rsp), %rax
	movq	%rax, 488(%rsp)
	leaq	80(%rsp), %rax
	movq	%rax, 496(%rsp)
	leaq	72(%rsp), %rax
	movq	%rax, 504(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 512(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 520(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 528(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 536(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 544(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 552(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 560(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 568(%rsp)
	movq	%rsp, %rax
	movq	%rax, 576(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 272(%rsp)
	movl	$1, 280(%rsp)
	movq	%rax, 284(%rsp)
	movl	$1, 292(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L6
	pushq	264(%rsp)
	.cfi_def_cfa_offset 616
	leaq	_Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi(%rip), %rdi
	pushq	264(%rsp)
	.cfi_def_cfa_offset 624
	movq	300(%rsp), %rcx
	movl	308(%rsp), %r8d
	movq	288(%rsp), %rsi
	movl	296(%rsp), %edx
	leaq	336(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 616
	popq	%rdx
	.cfi_def_cfa_offset 608
.L6:
	movq	584(%rsp), %rax
	subq	%fs:40, %rax
	je	.L8
	call	__stack_chk_fail@PLT
.L8:
	addq	$600, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2034:
	.size	_Z124__device_stub__Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_PiiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi, .-_Z124__device_stub__Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_PiiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi
	.globl	_Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi
	.type	_Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi, @function
_Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi:
.LFB2035:
	.cfi_startproc
	endbr64
	jmp	_Z124__device_stub__Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_PiiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi
	.cfi_endproc
.LFE2035:
	.size	_Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi, .-_Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2037:
	.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	.LC0(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi(%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
.LFE2037:
	.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 <hip/hip_runtime.h>
__device__ int find_qmin_and_qmax(
        double dq0, 
        double dq1, 
        double dq2, 
        double *qmin, 
        double *qmax)
{
    // Considering the centroid of an FV triangle and the vertices of its 
    // auxiliary triangle, find 
    // qmin=min(q)-qc and qmax=max(q)-qc, 
    // where min(q) and max(q) are respectively min and max over the
    // four values (at the centroid of the FV triangle and the auxiliary 
    // triangle vertices),
    // and qc is the centroid
    // dq0=q(vertex0)-q(centroid of FV triangle)
    // dq1=q(vertex1)-q(vertex0)
    // dq2=q(vertex2)-q(vertex0)

    // This is a simple implementation 
    *qmax = max(max(dq0, max(dq0+dq1, dq0+dq2)), 0.0) ;
    *qmin = min(min(dq0, min(dq0+dq1, dq0+dq2)), 0.0) ;

    return 0;
}



__device__ int limit_gradient(
        double *dqv, 
        double qmin, 
        double qmax, 
        double beta_w)
{
    // Given provisional jumps dqv from the FV triangle centroid to its 
    // vertices and jumps qmin (qmax) between the centroid of the FV 
    // triangle and the minimum (maximum) of the values at the centroid of 
    // the FV triangle and the auxiliary triangle vertices,
    // calculate a multiplicative factor phi by which the provisional 
    // vertex jumps are to be limited

    int i;
    double r=1000.0, r0=1.0, phi=1.0;
    double TINY = 1.0e-100; // to avoid machine accuracy problems.
    // FIXME: Perhaps use the epsilon used elsewhere.

    // Any provisional jump with magnitude < TINY does not contribute to 
    // the limiting process.

    for (i=0;i<3;i++){
        if (dqv[i]<-TINY)
            r0=qmin/dqv[i];

        if (dqv[i]>TINY)
            r0=qmax/dqv[i];

        r=min(r0,r);
    }

    phi=min(r*beta_w,1.0);
    //phi=1.;
    dqv[0]=dqv[0]*phi;
    dqv[1]=dqv[1]*phi;
    dqv[2]=dqv[2]*phi;

    return 0;
}



// Computational routine
__global__ void  _extrapolate_second_order_edge_sw(
        //int number_of_elements,
        int N,
        int optimise_dry_cells, 
        int extrapolate_velocity_second_order,
        double epsilon,
        double minimum_allowed_height,
        double beta_w,
        double beta_w_dry,
        double beta_uh,
        double beta_uh_dry,
        double beta_vh,
        double beta_vh_dry,

        long* surrogate_neighbours,
        long* number_of_boundaries,
        double* centroid_coordinates,

        double* stage_centroid_values,
        double* elevation_centroid_values,
        double* xmom_centroid_values,
        double* ymom_centroid_values,

        double* edge_coordinates,

        double* stage_edge_values,
        double* elevation_edge_values,
        double* xmom_edge_values,
        double* ymom_edge_values,

        double* stage_vertex_values,
        double* elevation_vertex_values,
        double* xmom_vertex_values,
        double* ymom_vertex_values,

        double * stage_centroid_store,
        double * xmom_centroid_store,
        double * ymom_centroid_store,
        double * min_elevation_edgevalue,
        double * max_elevation_edgevalue,
        int * count_wet_neighbours
        ) 
{
    const int k = 
        threadIdx.x+threadIdx.y*blockDim.x+
        (blockIdx.x+blockIdx.y*gridDim.x)*blockDim.x*blockDim.y;

    if ( k >= N)
        return;
    // Local variables
    // Gradient vector used to calculate edge values from centroids
    double a, b; 
    int k0, k1, k2, k3, k6, coord_index, i, ii, ktmp;
    // Vertices of the auxiliary triangle
    double x, y, x0, y0, x1, y1, x2, y2, xv0, yv0, xv1, yv1, xv2, yv2; 
    double dx1, dx2, dy1, dy2;
    double dxv0, dxv1, dxv2, dyv0, dyv1, dyv2, dq0, dq1, dq2;
    double area2, inv_area2;
    double dqv[3], qmin, qmax, hmin, hmax, bedmax, stagemin;
    double hc, h0, h1, h2, beta_tmp, hfactor, xtmp, ytmp;
    double dk, dv0, dv1, dv2, de[3];
    double demin, dcmax, r0scale, vect_norm, l1, l2;

    //double *xmom_centroid_store, *ymom_centroid_store; 
    //double *stage_centroid_store, *min_elevation_edgevalue;
    //double *max_elevation_edgevalue;
    //int *count_wet_neighbours;

    // Use malloc to avoid putting these variables on the stack, 
    // which can cause
    // segfaults in large model runs
    //xmom_centroid_store = malloc(number_of_elements*sizeof(double));
    //ymom_centroid_store = malloc(number_of_elements*sizeof(double));
    //stage_centroid_store = malloc(number_of_elements*sizeof(double));
    //min_elevation_edgevalue = malloc(number_of_elements*sizeof(double));
    //max_elevation_edgevalue = malloc(number_of_elements*sizeof(double));
    //count_wet_neighbours = malloc(number_of_elements*sizeof(int));

    if(extrapolate_velocity_second_order==1){
        // Replace momentum centroid with velocity centroid 2 allow velocity
        // extrapolation This will be changed back at the end of the routine
        //for (k=0; k<number_of_elements; k++){

        dk = max(stage_centroid_values[k]-elevation_centroid_values[k],
                minimum_allowed_height);
        xmom_centroid_store[k] = xmom_centroid_values[k];
        xmom_centroid_values[k] = xmom_centroid_values[k]/dk;

        ymom_centroid_store[k] = ymom_centroid_values[k];
        ymom_centroid_values[k] = ymom_centroid_values[k]/dk;

        min_elevation_edgevalue[k] = min(elevation_edge_values[3*k], 
                min(elevation_edge_values[3*k+1],
                    elevation_edge_values[3*k+2]));
        max_elevation_edgevalue[k] = max(elevation_edge_values[3*k], 
                max(elevation_edge_values[3*k+1],
                    elevation_edge_values[3*k+2]));
        //} // for k=0 to number_of_elements-1
    }

    // Count how many 'fully submerged' neighbours the cell has
    //for(k=0; k<number_of_elements;k++){ 
    count_wet_neighbours[k]=0;
    for (i=0; i<3; i++){
        ktmp = surrogate_neighbours[3*k+i];              
        if(stage_centroid_values[ktmp] > max_elevation_edgevalue[ktmp]){
            count_wet_neighbours[k]+=1;
        }
    }
    //} // for k=0 to number_of_elements-1

    // Begin extrapolation routine
    //for (k = 0; k < number_of_elements; k++) {
    do{
    k3=k*3;
    k6=k*6;

    if (number_of_boundaries[k]==3)
        //if (0==0)
    {
        // No neighbours, set gradient on the triangle to zero

        stage_edge_values[k3]   = stage_centroid_values[k];
        stage_edge_values[k3+1] = stage_centroid_values[k];
        stage_edge_values[k3+2] = stage_centroid_values[k];
        xmom_edge_values[k3]    = xmom_centroid_values[k];
        xmom_edge_values[k3+1]  = xmom_centroid_values[k];
        xmom_edge_values[k3+2]  = xmom_centroid_values[k];
        ymom_edge_values[k3]    = ymom_centroid_values[k];
        ymom_edge_values[k3+1]  = ymom_centroid_values[k];
        ymom_edge_values[k3+2]  = ymom_centroid_values[k];

        continue;
    }
    else 
    {
        // Triangle k has one or more neighbours. 
        // Get centroid and edge coordinates of the triangle

        // Get the edge coordinates
        xv0 = edge_coordinates[k6];   
        yv0 = edge_coordinates[k6+1];
        xv1 = edge_coordinates[k6+2]; 
        yv1 = edge_coordinates[k6+3];
        xv2 = edge_coordinates[k6+4]; 
        yv2 = edge_coordinates[k6+5];

        // Get the centroid coordinates
        coord_index = 2*k;
        x = centroid_coordinates[coord_index];
        y = centroid_coordinates[coord_index+1];

        // Store x- and y- differentials for the edges of 
        // triangle k relative to the centroid
        dxv0 = xv0 - x; 
        dxv1 = xv1 - x; 
        dxv2 = xv2 - x;
        dyv0 = yv0 - y; 
        dyv1 = yv1 - y; 
        dyv2 = yv2 - y;
        // Compute the minimum distance from the centroid to an edge
        //demin=min(dxv0*dxv0 +dyv0*dyv0, 
        //      min(dxv1*dxv1+dyv1*dyv1, dxv2*dxv2+dyv2*dyv2));
        //demin=sqrt(demin);
    }



    if (number_of_boundaries[k]<=1)
    {
        //==============================================
        // Number of boundaries <= 1
        //==============================================    


        // If no boundaries, auxiliary triangle is formed 
        // from the centroids of the three neighbours
        // If one boundary, auxiliary triangle is formed 
        // from this centroid and its two neighbours

        k0 = surrogate_neighbours[k3];
        k1 = surrogate_neighbours[k3 + 1];
        k2 = surrogate_neighbours[k3 + 2];

        // Take note if the max neighbour bed elevation 
        // is greater than the min
        // neighbour stage -- suggests a 'steep' bed relative to the flow
        bedmax = max(elevation_centroid_values[k], 
                    max(elevation_centroid_values[k0],
                        max(elevation_centroid_values[k1], 
                            elevation_centroid_values[k2])));
        //bedmax = elevation_centroid_values[k];
        stagemin = min(max(stage_centroid_values[k],
                            elevation_centroid_values[k]), 
                min(max(stage_centroid_values[k0], 
                        elevation_centroid_values[k0]),
                    min(max(stage_centroid_values[k1], 
                        elevation_centroid_values[k1]),
                        max(stage_centroid_values[k2], 
                            elevation_centroid_values[k2]))));
        if(stagemin < bedmax){
            // This will cause first order extrapolation
            k2 = k;
            k0 = k;
            k1 = k;
        } 

        // Get the auxiliary triangle's vertex coordinates 
        // (really the centroids of neighbouring triangles)
        coord_index = 2*k0;
        x0 = centroid_coordinates[coord_index];
        y0 = centroid_coordinates[coord_index+1];

        coord_index = 2*k1;
        x1 = centroid_coordinates[coord_index];
        y1 = centroid_coordinates[coord_index+1];

        coord_index = 2*k2;
        x2 = centroid_coordinates[coord_index];
        y2 = centroid_coordinates[coord_index+1];

        // Store x- and y- differentials for the vertices 
        // of the auxiliary triangle
        dx1 = x1 - x0; 
        dx2 = x2 - x0;
        dy1 = y1 - y0; 
        dy2 = y2 - y0;

        // Calculate 2*area of the auxiliary triangle
        // The triangle is guaranteed to be counter-clockwise      
        area2 = dy2*dx1 - dy1*dx2; 

        // Treat triangles with zero or 1 wet neighbours. 
        if ((area2 <= 0)) //|(count_wet_neighbours[k]==0))
        {
            //printf("Error negative triangle area \n");
            //report_python_error(AT, "Negative triangle area");
            //return -1;
            stage_edge_values[k3]   = stage_centroid_values[k];
            stage_edge_values[k3+1] = stage_centroid_values[k];
            stage_edge_values[k3+2] = stage_centroid_values[k];
            // First order momentum / velocity extrapolation
            xmom_edge_values[k3]    = xmom_centroid_values[k];
            xmom_edge_values[k3+1]  = xmom_centroid_values[k];
            xmom_edge_values[k3+2]  = xmom_centroid_values[k];
            ymom_edge_values[k3]    = ymom_centroid_values[k];
            ymom_edge_values[k3+1]  = ymom_centroid_values[k];
            ymom_edge_values[k3+2]  = ymom_centroid_values[k];

            continue;
        }  

        // Calculate heights of neighbouring cells
        hc = stage_centroid_values[k]  - elevation_centroid_values[k];
        h0 = stage_centroid_values[k0] - elevation_centroid_values[k0];
        h1 = stage_centroid_values[k1] - elevation_centroid_values[k1];
        h2 = stage_centroid_values[k2] - elevation_centroid_values[k2];
        hmin = min(min(h0, min(h1, h2)), hc);

        hfactor = 0.0;
        //if (hmin > 0.001) 
        if (hmin > 0.) 
            //if (hc>0.0)
        {
            hfactor = 1.0 ;//hmin/(hmin + 0.004);
            //hfactor=hmin/(hmin + 0.004);
        }


        //-----------------------------------
        // stage
        //-----------------------------------      

        // Calculate the difference between vertex 0 of the auxiliary 
        // triangle and the centroid of triangle k
        dq0 = stage_centroid_values[k0] - stage_centroid_values[k];

        // Calculate differentials between the vertices 
        // of the auxiliary triangle (centroids of neighbouring triangles)
        dq1 = stage_centroid_values[k1] - stage_centroid_values[k0];
        dq2 = stage_centroid_values[k2] - stage_centroid_values[k0];

        inv_area2 = 1.0/area2;
        // Calculate the gradient of stage on the auxiliary triangle
        a = dy2*dq1 - dy1*dq2;
        a *= inv_area2;
        b = dx1*dq2 - dx2*dq1;
        b *= inv_area2;

        // Calculate provisional jumps in stage from the centroid 
        // of triangle k to its vertices, to be limited
        dqv[0] = a*dxv0 + b*dyv0;
        dqv[1] = a*dxv1 + b*dyv1;
        dqv[2] = a*dxv2 + b*dyv2;

        // Now we want to find min and max of the centroid and the 
        // vertices of the auxiliary triangle and compute jumps 
        // from the centroid to the min and max
        find_qmin_and_qmax(dq0, dq1, dq2, &qmin, &qmax);

        beta_tmp = beta_w_dry + (beta_w - beta_w_dry) * hfactor;


        // Limit the gradient
        limit_gradient(dqv, qmin, qmax, beta_tmp);
        stage_edge_values[k3+0] = stage_centroid_values[k] + dqv[0];
        stage_edge_values[k3+1] = stage_centroid_values[k] + dqv[1];
        stage_edge_values[k3+2] = stage_centroid_values[k] + dqv[2];


        //-----------------------------------
        // xmomentum
        //-----------------------------------            

        // Calculate the difference between vertex 0 of the auxiliary 
        // triangle and the centroid of triangle k      
        dq0 = xmom_centroid_values[k0] - xmom_centroid_values[k];

        // Calculate differentials between the vertices 
        // of the auxiliary triangle
        dq1 = xmom_centroid_values[k1] - xmom_centroid_values[k0];
        dq2 = xmom_centroid_values[k2] - xmom_centroid_values[k0];

        // Calculate the gradient of xmom on the auxiliary triangle
        a = dy2*dq1 - dy1*dq2;
        a *= inv_area2;
        b = dx1*dq2 - dx2*dq1;
        b *= inv_area2;

        // Calculate provisional jumps in stage from the centroid 
        // of triangle k to its vertices, to be limited      
        dqv[0] = a*dxv0+b*dyv0;
        dqv[1] = a*dxv1+b*dyv1;
        dqv[2] = a*dxv2+b*dyv2;

        // Now we want to find min and max of the centroid and the 
        // vertices of the auxiliary triangle and compute jumps 
        // from the centroid to the min and max
        //
        find_qmin_and_qmax(dq0, dq1, dq2, &qmin, &qmax);

        beta_tmp = beta_uh_dry + (beta_uh - beta_uh_dry) * hfactor;

        limit_gradient(dqv, qmin, qmax, beta_tmp);


        for (i=0; i < 3; i++)
        {
            xmom_edge_values[k3+i] = xmom_centroid_values[k] + dqv[i];
        }

        //-----------------------------------
        // ymomentum
        //-----------------------------------                  

        // Calculate the difference between vertex 0 of the auxiliary 
        // triangle and the centroid of triangle k
        dq0 = ymom_centroid_values[k0] - ymom_centroid_values[k];

        // Calculate differentials between the vertices 
        // of the auxiliary triangle
        dq1 = ymom_centroid_values[k1] - ymom_centroid_values[k0];
        dq2 = ymom_centroid_values[k2] - ymom_centroid_values[k0];

        // Calculate the gradient of xmom on the auxiliary triangle
        a = dy2*dq1 - dy1*dq2;
        a *= inv_area2;
        b = dx1*dq2 - dx2*dq1;
        b *= inv_area2;

        // Calculate provisional jumps in stage from the centroid 
        // of triangle k to its vertices, to be limited
        dqv[0] = a*dxv0 + b*dyv0;
        dqv[1] = a*dxv1 + b*dyv1;
        dqv[2] = a*dxv2 + b*dyv2;

        // Now we want to find min and max of the centroid and the 
        // vertices of the auxiliary triangle and compute jumps 
        // from the centroid to the min and max
        //
        find_qmin_and_qmax(dq0, dq1, dq2, &qmin, &qmax);

        beta_tmp = beta_vh_dry + (beta_vh - beta_vh_dry) * hfactor;   

        limit_gradient(dqv, qmin, qmax, beta_tmp);

        for (i=0;i<3;i++)
        {
            ymom_edge_values[k3 + i] = ymom_centroid_values[k] + dqv[i];
        }

    } // End number_of_boundaries <=1 
    else
    {

        //==============================================
        // Number of boundaries == 2
        //==============================================        

        // One internal neighbour and gradient 
        // is in direction of the neighbour's centroid

        // Find the only internal neighbour (k1?)
        for (k2 = k3; k2 < k3 + 3; k2++)
        {
            // Find internal neighbour of triangle k      
            // k2 indexes the edges of triangle k   

            if (surrogate_neighbours[k2] != k)
            {
                break;
            }
        }

        if ((k2 == k3 + 3)) 
        {
            // If we didn't find an internal neighbour
            //report_python_error(AT, "Internal neighbour not found");
            //return -1;
            return;
        }

        k1 = surrogate_neighbours[k2];

        // The coordinates of the triangle are already (x,y). 
        // Get centroid of the neighbour (x1,y1)
        coord_index = 2*k1;
        x1 = centroid_coordinates[coord_index];
        y1 = centroid_coordinates[coord_index + 1];

        // Compute x- and y- distances between the centroid of 
        // triangle k and that of its neighbour
        dx1 = x1 - x; 
        dy1 = y1 - y;

        // Set area2 as the square of the distance
        area2 = dx1*dx1 + dy1*dy1;

        // Set dx2=(x1-x0)/((x1-x0)^2+(y1-y0)^2) 
        // and dy2=(y1-y0)/((x1-x0)^2+(y1-y0)^2) which
        // respectively correspond to the x- and y- gradients 
        // of the conserved quantities
        dx2 = 1.0/area2;
        dy2 = dx2*dy1;
        dx2 *= dx1;


        //-----------------------------------
        // stage
        //-----------------------------------            

        // Compute differentials
        dq1 = stage_centroid_values[k1] - stage_centroid_values[k];

        // Calculate the gradient between the centroid of triangle k 
        // and that of its neighbour
        a = dq1*dx2;
        b = dq1*dy2;

        // Calculate provisional edge jumps, to be limited
        dqv[0] = a*dxv0 + b*dyv0;
        dqv[1] = a*dxv1 + b*dyv1;
        dqv[2] = a*dxv2 + b*dyv2;

        // Now limit the jumps
        if (dq1>=0.0)
        {
            qmin=0.0;
            qmax=dq1;
        }
        else
        {
            qmin = dq1;
            qmax = 0.0;
        }

        // Limit the gradient
        limit_gradient(dqv, qmin, qmax, beta_w);

        //for (i=0; i < 3; i++)
        //{
        stage_edge_values[k3] = stage_centroid_values[k] + dqv[0];
        stage_edge_values[k3 + 1] = stage_centroid_values[k] + dqv[1];
        stage_edge_values[k3 + 2] = stage_centroid_values[k] + dqv[2];
        //}

        //-----------------------------------
        // xmomentum
        //-----------------------------------                        

        // Compute differentials
        dq1 = xmom_centroid_values[k1] - xmom_centroid_values[k];

        // Calculate the gradient between the centroid of triangle k 
        // and that of its neighbour
        a = dq1*dx2;
        b = dq1*dy2;

        // Calculate provisional edge jumps, to be limited
        dqv[0] = a*dxv0+b*dyv0;
        dqv[1] = a*dxv1+b*dyv1;
        dqv[2] = a*dxv2+b*dyv2;

        // Now limit the jumps
        if (dq1 >= 0.0)
        {
            qmin = 0.0;
            qmax = dq1;
        }
        else
        {
            qmin = dq1;
            qmax = 0.0;
        }

        // Limit the gradient      
        limit_gradient(dqv, qmin, qmax, beta_w);

        //for (i=0;i<3;i++)
        //xmom_edge_values[k3] = xmom_centroid_values[k] + dqv[0];
        //xmom_edge_values[k3 + 1] = xmom_centroid_values[k] + dqv[1];
        //xmom_edge_values[k3 + 2] = xmom_centroid_values[k] + dqv[2];

        for (i = 0; i < 3;i++)
        {
            xmom_edge_values[k3 + i] = xmom_centroid_values[k] + dqv[i];
        }

        //-----------------------------------
        // ymomentum
        //-----------------------------------                        

        // Compute differentials
        dq1 = ymom_centroid_values[k1] - ymom_centroid_values[k];

        // Calculate the gradient between the centroid of triangle k 
        // and that of its neighbours      a = dq1*dx2;
        b = dq1*dy2;

        // Calculate provisional edge jumps, to be limited
        dqv[0] = a*dxv0 + b*dyv0;
        dqv[1] = a*dxv1 + b*dyv1;
        dqv[2] = a*dxv2 + b*dyv2;

        // Now limit the jumps
        if (dq1>=0.0)
        {
            qmin = 0.0;
            qmax = dq1;
        } 
        else
        {
            qmin = dq1;
            qmax = 0.0;
        }

        // Limit the gradient
        limit_gradient(dqv, qmin, qmax, beta_w);

        for (i=0;i<3;i++)
        {
            ymom_edge_values[k3 + i] = ymom_centroid_values[k] + dqv[i];
        }
    } // else [number_of_boundaries==2]
    } while(k < 0);
    //} // for k=0 to number_of_elements-1

    // Compute vertex values of quantities
    //for (k=0; k<number_of_elements; k++){
    k3=3*k;

    // Compute stage vertex values 
    stage_vertex_values[k3] = stage_edge_values[k3+1] + stage_edge_values[k3+2] -stage_edge_values[k3] ; 
    stage_vertex_values[k3+1] =  stage_edge_values[k3] + stage_edge_values[k3+2]-stage_edge_values[k3+1]; 
    stage_vertex_values[k3+2] =  stage_edge_values[k3] + stage_edge_values[k3+1]-stage_edge_values[k3+2]; 

    // Compute xmom vertex values 
    xmom_vertex_values[k3] = xmom_edge_values[k3+1] + xmom_edge_values[k3+2] -xmom_edge_values[k3] ; 
    xmom_vertex_values[k3+1] =  xmom_edge_values[k3] + xmom_edge_values[k3+2]-xmom_edge_values[k3+1]; 
    xmom_vertex_values[k3+2] =  xmom_edge_values[k3] + xmom_edge_values[k3+1]-xmom_edge_values[k3+2]; 

    // Compute ymom vertex values 
    ymom_vertex_values[k3] = ymom_edge_values[k3+1] + ymom_edge_values[k3+2] -ymom_edge_values[k3] ; 
    ymom_vertex_values[k3+1] =  ymom_edge_values[k3] + ymom_edge_values[k3+2]-ymom_edge_values[k3+1]; 
    ymom_vertex_values[k3+2] =  ymom_edge_values[k3] + ymom_edge_values[k3+1]-ymom_edge_values[k3+2]; 

    // If needed, convert from velocity to momenta
    if(extrapolate_velocity_second_order==1){
        //Convert velocity back to momenta at centroids
        xmom_centroid_values[k] = xmom_centroid_store[k];
        ymom_centroid_values[k] = ymom_centroid_store[k];

        // Re-compute momenta at edges
        for (i=0; i<3; i++){
            de[i] = max(stage_edge_values[k3+i]-elevation_edge_values[k3+i],0.0);
            xmom_edge_values[k3+i]=xmom_edge_values[k3+i]*de[i];
            ymom_edge_values[k3+i]=ymom_edge_values[k3+i]*de[i];
        }

        // Re-compute momenta at vertices
        for (i=0; i<3; i++){
            de[i] = max(stage_vertex_values[k3+i]-elevation_vertex_values[k3+i],0.0);
            xmom_vertex_values[k3+i]=xmom_vertex_values[k3+i]*de[i];
            ymom_vertex_values[k3+i]=ymom_vertex_values[k3+i]*de[i];
        }

    }
    //} // for k=0 to number_of_elements-1

}  

	.text
	.file	"extrapolate_second_order_edge_swb2.hip"
	.globl	_Z48__device_stub___extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi # -- Begin function _Z48__device_stub___extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi
	.type	_Z48__device_stub___extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi,@function
_Z48__device_stub___extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi: # @_Z48__device_stub___extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi
	.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	$424, %rsp                      # imm = 0x1A8
	.cfi_def_cfa_offset 480
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	leaq	12(%rsp), %rax
	movl	%edi, (%rax)
	leaq	20(%rsp), %rdi
	movl	%esi, (%rdi)
	leaq	16(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	120(%rsp), %rdx
	movsd	%xmm0, (%rdx)
	leaq	112(%rsp), %r10
	movsd	%xmm1, (%r10)
	leaq	104(%rsp), %r11
	movsd	%xmm2, (%r11)
	leaq	96(%rsp), %r14
	movsd	%xmm3, (%r14)
	leaq	88(%rsp), %r15
	movsd	%xmm4, (%r15)
	leaq	80(%rsp), %r12
	movsd	%xmm5, (%r12)
	leaq	72(%rsp), %r13
	movsd	%xmm6, (%r13)
	leaq	64(%rsp), %rbp
	movsd	%xmm7, (%rbp)
	leaq	56(%rsp), %rax
	movq	%rcx, (%rax)
	leaq	48(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	40(%rsp), %r8
	movq	%r9, (%r8)
	leaq	160(%rsp), %rbx
	leaq	12(%rsp), %r9
	movq	%r9, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%r10, 32(%rbx)
	movq	%r11, 40(%rbx)
	movq	%r14, 48(%rbx)
	movq	%r15, 56(%rbx)
	movq	%r12, 64(%rbx)
	movq	%r13, 72(%rbx)
	movq	%rbp, 80(%rbx)
	movq	%rax, 88(%rbx)
	movq	%rcx, 96(%rbx)
	movq	%r8, 104(%rbx)
	leaq	480(%rsp), %rax
	movq	%rax, 112(%rbx)
	leaq	488(%rsp), %rax
	movq	%rax, 120(%rbx)
	leaq	496(%rsp), %rax
	movq	%rax, 128(%rbx)
	leaq	504(%rsp), %rax
	movq	%rax, 136(%rbx)
	leaq	512(%rsp), %rax
	movq	%rax, 144(%rbx)
	leaq	520(%rsp), %rax
	movq	%rax, 152(%rbx)
	leaq	528(%rsp), %rax
	movq	%rax, 160(%rbx)
	leaq	536(%rsp), %rax
	movq	%rax, 168(%rbx)
	leaq	544(%rsp), %rax
	movq	%rax, 176(%rbx)
	leaq	552(%rsp), %rax
	movq	%rax, 184(%rbx)
	leaq	560(%rsp), %rax
	movq	%rax, 192(%rbx)
	leaq	568(%rsp), %rax
	movq	%rax, 200(%rbx)
	leaq	576(%rsp), %rax
	movq	%rax, 208(%rbx)
	leaq	584(%rsp), %rax
	movq	%rax, 216(%rbx)
	leaq	592(%rsp), %rax
	movq	%rax, 224(%rbx)
	leaq	600(%rsp), %rax
	movq	%rax, 232(%rbx)
	leaq	608(%rsp), %rax
	movq	%rax, 240(%rbx)
	leaq	616(%rsp), %rax
	movq	%rax, 248(%rbx)
	leaq	624(%rsp), %rax
	movq	%rax, 256(%rbx)
	leaq	144(%rsp), %r14
	leaq	128(%rsp), %r15
	leaq	32(%rsp), %r12
	leaq	24(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$440, %rsp                      # imm = 0x1B8
	.cfi_adjust_cfa_offset -440
	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_end0:
	.size	_Z48__device_stub___extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi, .Lfunc_end0-_Z48__device_stub___extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi
	.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	.LBB1_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB1_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi, %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_end1:
	.size	__hip_module_ctor, .Lfunc_end1-__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	.LBB2_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
.LBB2_2:
	retq
.Lfunc_end2:
	.size	__hip_module_dtor, .Lfunc_end2-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi,@object # @_Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi
	.section	.rodata,"a",@progbits
	.globl	_Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi
	.p2align	3, 0x0
_Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi:
	.quad	_Z48__device_stub___extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi
	.size	_Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi, 8

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

	.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 _Z48__device_stub___extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi ; -- Begin function _Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi
	.globl	_Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi
	.p2align	8
	.type	_Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi,@function
_Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi: ; @_Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi
; %bb.0:
	s_clause 0x2
	s_load_b32 s2, s[0:1], 0x100
	s_load_b32 s3, s[0:1], 0x10c
	s_load_b32 s4, s[0:1], 0x0
	v_bfe_u32 v1, v0, 10, 10
	v_and_b32_e32 v0, 0x3ff, v0
	s_waitcnt lgkmcnt(0)
	s_mul_i32 s2, s2, s15
	s_lshr_b32 s5, s3, 16
	s_add_i32 s2, s2, s14
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mad_u64_u32 v[2:3], null, s2, s5, v[1:2]
	s_and_b32 s2, s3, 0xffff
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_mad_u64_u32 v[6:7], null, v2, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_i32_e64 s4, v6
	s_cbranch_execz .LBB0_60
; %bb.1:
	s_clause 0x6
	s_load_b32 s33, s[0:1], 0x8
	s_load_b512 s[4:19], s[0:1], 0x58
	s_load_b64 s[34:35], s[0:1], 0xc8
	s_load_b128 s[28:31], s[0:1], 0xb8
	s_load_b256 s[20:27], s[0:1], 0x98
	s_load_b64 s[2:3], s[0:1], 0xf8
	s_load_b256 s[52:59], s[0:1], 0xd8
	v_ashrrev_i32_e32 v7, 31, v6
	v_lshl_add_u32 v0, v6, 1, v6
	s_waitcnt lgkmcnt(0)
	s_cmp_eq_u32 s33, 1
	s_cselect_b32 s33, -1, 0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 vcc_lo, exec_lo, s33
	s_cbranch_vccnz .LBB0_3
; %bb.2:                                ; %._crit_edge
	v_lshl_add_u32 v8, v6, 1, v6
	s_mov_b32 s60, 0
	s_branch .LBB0_4
.LBB0_3:
	s_mov_b32 s60, -1
                                        ; implicit-def: $vgpr8
.LBB0_4:                                ; %Flow1028
	s_load_b512 s[36:51], s[0:1], 0x18
	v_lshlrev_b64 v[4:5], 3, v[6:7]
	s_and_not1_b32 vcc_lo, exec_lo, s60
	s_cbranch_vccnz .LBB0_6
; %bb.5:
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v1, vcc_lo, s8, v4
	v_add_co_ci_u32_e32 v2, vcc_lo, s9, v5, vcc_lo
	v_add_co_u32 v8, vcc_lo, s10, v4
	v_add_co_ci_u32_e32 v9, vcc_lo, s11, v5, vcc_lo
	v_add_co_u32 v10, vcc_lo, s12, v4
	global_load_b64 v[1:2], v[1:2], off
	global_load_b64 v[8:9], v[8:9], off
	v_add_co_ci_u32_e32 v11, vcc_lo, s13, v5, vcc_lo
	s_waitcnt lgkmcnt(0)
	v_max_f64 v[14:15], s[36:37], s[36:37]
	global_load_b64 v[12:13], v[10:11], off
	s_waitcnt vmcnt(1)
	v_add_f64 v[1:2], v[1:2], -v[8:9]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_max_f64 v[1:2], v[1:2], v[14:15]
	s_waitcnt vmcnt(0)
	v_div_scale_f64 v[8:9], null, v[1:2], v[1:2], v[12:13]
	v_div_scale_f64 v[18:19], vcc_lo, v[12:13], v[1:2], v[12:13]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f64_e32 v[14:15], v[8:9]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[16:17], -v[8:9], v[14:15], 1.0
	v_fma_f64 v[14:15], v[14:15], v[16:17], v[14:15]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[16:17], -v[8:9], v[14:15], 1.0
	v_fma_f64 v[14:15], v[14:15], v[16:17], v[14:15]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[16:17], v[18:19], v[14:15]
	v_fma_f64 v[8:9], -v[8:9], v[16:17], v[18:19]
	s_delay_alu instid0(VALU_DEP_1)
	v_div_fmas_f64 v[8:9], v[8:9], v[14:15], v[16:17]
	v_add_co_u32 v14, vcc_lo, s52, v4
	v_add_co_ci_u32_e32 v15, vcc_lo, s53, v5, vcc_lo
	v_add_co_u32 v16, vcc_lo, s14, v4
	v_add_co_ci_u32_e32 v17, vcc_lo, s15, v5, vcc_lo
	v_div_fixup_f64 v[8:9], v[8:9], v[1:2], v[12:13]
	global_store_b64 v[14:15], v[12:13], off
	global_store_b64 v[10:11], v[8:9], off
	global_load_b64 v[8:9], v[16:17], off
	s_waitcnt vmcnt(0)
	v_div_scale_f64 v[10:11], null, v[1:2], v[1:2], v[8:9]
	v_div_scale_f64 v[18:19], vcc_lo, v[8:9], v[1:2], v[8:9]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f64_e32 v[12:13], v[10:11]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[14:15], -v[10:11], v[12:13], 1.0
	v_fma_f64 v[12:13], v[12:13], v[14:15], v[12:13]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[14:15], -v[10:11], v[12:13], 1.0
	v_fma_f64 v[12:13], v[12:13], v[14:15], v[12:13]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[14:15], v[18:19], v[12:13]
	v_fma_f64 v[10:11], -v[10:11], v[14:15], v[18:19]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_div_fmas_f64 v[10:11], v[10:11], v[12:13], v[14:15]
	v_add_co_u32 v12, vcc_lo, s54, v4
	v_add_co_ci_u32_e32 v13, vcc_lo, s55, v5, vcc_lo
	v_div_fixup_f64 v[2:3], v[10:11], v[1:2], v[8:9]
	v_ashrrev_i32_e32 v1, 31, v0
	global_store_b64 v[12:13], v[8:9], off
	global_store_b64 v[16:17], v[2:3], off
	v_lshlrev_b64 v[10:11], 3, v[0:1]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v14, vcc_lo, s20, v10
	v_add_co_ci_u32_e32 v15, vcc_lo, s21, v11, vcc_lo
	s_clause 0x1
	global_load_b64 v[1:2], v[14:15], off offset:16
	global_load_b128 v[8:11], v[14:15], off
	s_waitcnt vmcnt(1)
	v_max_f64 v[1:2], v[1:2], v[1:2]
	s_waitcnt vmcnt(0)
	v_max_f64 v[10:11], v[10:11], v[10:11]
	v_max_f64 v[8:9], v[8:9], v[8:9]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_min_f64 v[1:2], v[10:11], v[1:2]
	v_min_f64 v[1:2], v[8:9], v[1:2]
	v_add_co_u32 v8, vcc_lo, s56, v4
	v_add_co_ci_u32_e32 v9, vcc_lo, s57, v5, vcc_lo
	global_store_b64 v[8:9], v[1:2], off
	s_clause 0x1
	global_load_b64 v[1:2], v[14:15], off offset:16
	global_load_b128 v[8:11], v[14:15], off
	s_waitcnt vmcnt(1)
	v_max_f64 v[1:2], v[1:2], v[1:2]
	s_waitcnt vmcnt(0)
	v_max_f64 v[10:11], v[10:11], v[10:11]
	v_max_f64 v[8:9], v[8:9], v[8:9]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_max_f64 v[1:2], v[10:11], v[1:2]
	v_max_f64 v[1:2], v[8:9], v[1:2]
	v_add_co_u32 v9, vcc_lo, s58, v4
	v_add_co_ci_u32_e32 v10, vcc_lo, s59, v5, vcc_lo
	v_mov_b32_e32 v8, v0
	global_store_b64 v[9:10], v[1:2], off
.LBB0_6:
	v_ashrrev_i32_e32 v9, 31, v8
	v_lshlrev_b64 v[2:3], 2, v[6:7]
	v_mov_b32_e32 v10, 0
	s_mov_b64 s[0:1], 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[0:1], 3, v[8:9]
	v_add_co_u32 v2, vcc_lo, s2, v2
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v3, vcc_lo, s3, v3, vcc_lo
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v14, vcc_lo, s50, v0
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_ci_u32_e32 v15, vcc_lo, s51, v1, vcc_lo
	global_store_b32 v[2:3], v10, off
.LBB0_7:                                ; =>This Inner Loop Header: Depth=1
	v_add_co_u32 v11, vcc_lo, v14, s0
	v_add_co_ci_u32_e32 v12, vcc_lo, s1, v15, vcc_lo
	s_mov_b32 s2, exec_lo
	global_load_b32 v11, v[11:12], off
	s_waitcnt vmcnt(0)
	v_ashrrev_i32_e32 v12, 31, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[11:12], 3, v[11:12]
	v_add_co_u32 v16, vcc_lo, s8, v11
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v17, vcc_lo, s9, v12, vcc_lo
	v_add_co_u32 v11, vcc_lo, s58, v11
	v_add_co_ci_u32_e32 v12, vcc_lo, s59, v12, vcc_lo
	global_load_b64 v[16:17], v[16:17], off
	global_load_b64 v[11:12], v[11:12], off
	s_waitcnt vmcnt(0)
	v_cmpx_gt_f64_e32 v[16:17], v[11:12]
	s_cbranch_execz .LBB0_9
; %bb.8:                                ;   in Loop: Header=BB0_7 Depth=1
	v_add_nc_u32_e32 v10, 1, v10
	global_store_b32 v[2:3], v10, off
.LBB0_9:                                ;   in Loop: Header=BB0_7 Depth=1
	s_or_b32 exec_lo, exec_lo, s2
	s_add_u32 s0, s0, 8
	s_addc_u32 s1, s1, 0
	s_cmp_eq_u32 s0, 24
	s_cbranch_scc0 .LBB0_7
; %bb.10:                               ; %.preheader769
	v_add_co_u32 v2, vcc_lo, s4, v4
	v_add_co_ci_u32_e32 v3, vcc_lo, s5, v5, vcc_lo
	v_add_f64 v[21:22], s[38:39], -s[40:41]
	v_add_f64 v[23:24], s[42:43], -s[44:45]
	v_add_f64 v[25:26], s[46:47], -s[48:49]
	global_load_b64 v[2:3], v[2:3], off
	v_add_co_u32 v17, vcc_lo, s8, v4
	v_add_co_ci_u32_e32 v18, vcc_lo, s9, v5, vcc_lo
	v_add_co_u32 v19, vcc_lo, s18, v0
	v_add_co_ci_u32_e32 v20, vcc_lo, s19, v1, vcc_lo
	v_add_co_u32 v12, vcc_lo, s12, v4
	v_add_co_ci_u32_e32 v13, vcc_lo, s13, v5, vcc_lo
	v_add_co_u32 v10, vcc_lo, s14, v4
	v_mul_lo_u32 v35, v6, 6
	v_lshlrev_b32_e32 v33, 1, v6
	v_add_co_ci_u32_e32 v11, vcc_lo, s15, v5, vcc_lo
	v_add_co_u32 v27, vcc_lo, s50, v0
	v_add_co_ci_u32_e32 v28, vcc_lo, s51, v1, vcc_lo
	v_add_co_u32 v29, vcc_lo, s10, v4
	v_ashrrev_i32_e32 v34, 31, v33
	v_add_co_ci_u32_e32 v30, vcc_lo, s11, v5, vcc_lo
	v_add_co_u32 v31, vcc_lo, s22, v0
	v_ashrrev_i32_e32 v36, 31, v35
	v_add_co_ci_u32_e32 v32, vcc_lo, s23, v1, vcc_lo
	v_lshlrev_b64 v[37:38], 3, v[33:34]
	v_add_co_u32 v33, vcc_lo, s24, v0
	v_add_co_ci_u32_e32 v34, vcc_lo, s25, v1, vcc_lo
	v_lshlrev_b64 v[0:1], 3, v[35:36]
	v_add_nc_u32_e32 v9, 3, v8
	v_dual_mov_b32 v39, 0 :: v_dual_add_nc_u32 v16, 1, v8
	v_add_co_u32 v35, vcc_lo, s6, v37
	v_add_co_ci_u32_e32 v36, vcc_lo, s7, v38, vcc_lo
	v_add_co_u32 v37, vcc_lo, s16, v0
	v_cmp_lt_i32_e64 s0, -1, v6
	v_max_i32_e32 v95, v9, v16
	v_add_co_ci_u32_e32 v38, vcc_lo, s17, v1, vcc_lo
	s_mov_b32 s4, 0xe48e0530
	s_mov_b32 s5, 0xab2bff2e
	s_mov_b32 s42, 0
	s_mov_b32 s17, 0x2b2bff2e
                                        ; implicit-def: $sgpr43
                                        ; implicit-def: $sgpr46
	s_waitcnt vmcnt(0)
	v_cmp_ne_u64_e64 s1, 3, v[2:3]
	v_cmp_lt_i64_e64 s2, 1, v[2:3]
.LBB0_11:                               ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_14 Depth 2
                                        ;     Child Loop BB0_19 Depth 2
                                        ;     Child Loop BB0_21 Depth 2
                                        ;     Child Loop BB0_23 Depth 2
                                        ;     Child Loop BB0_25 Depth 2
                                        ;     Child Loop BB0_27 Depth 2
                                        ;     Child Loop BB0_35 Depth 2
                                        ;     Child Loop BB0_37 Depth 2
                                        ;     Child Loop BB0_39 Depth 2
                                        ;     Child Loop BB0_41 Depth 2
                                        ;     Child Loop BB0_43 Depth 2
	s_mov_b32 s16, 0
	s_mov_b32 s3, 0
                                        ; implicit-def: $sgpr56
                                        ; implicit-def: $vgpr58_vgpr59
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_saveexec_b32 s36, s1
	s_xor_b32 s47, exec_lo, s36
	s_cbranch_execz .LBB0_47
; %bb.12:                               ;   in Loop: Header=BB0_11 Depth=1
	global_load_b128 v[40:43], v[37:38], off
	global_load_b128 v[0:3], v[35:36], off
	s_clause 0x1
	global_load_b128 v[51:54], v[37:38], off offset:16
	global_load_b128 v[55:58], v[37:38], off offset:32
	s_mov_b32 s59, 0
	s_mov_b32 s57, 0
                                        ; implicit-def: $sgpr56
	s_waitcnt vmcnt(2)
	v_add_f64 v[47:48], v[40:41], -v[0:1]
	s_waitcnt vmcnt(1)
	v_add_f64 v[49:50], v[51:52], -v[0:1]
	s_waitcnt vmcnt(0)
	v_add_f64 v[51:52], v[55:56], -v[0:1]
	v_add_f64 v[41:42], v[42:43], -v[2:3]
	v_add_f64 v[43:44], v[53:54], -v[2:3]
	v_add_f64 v[45:46], v[57:58], -v[2:3]
	s_and_saveexec_b32 s3, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s58, exec_lo, s3
	s_cbranch_execz .LBB0_30
; %bb.13:                               ; %.preheader768.preheader
                                        ;   in Loop: Header=BB0_11 Depth=1
	v_dual_mov_b32 v54, v15 :: v_dual_mov_b32 v53, v14
	v_mov_b32_e32 v56, v8
                                        ; implicit-def: $sgpr36
.LBB0_14:                               ; %.preheader768
                                        ;   Parent Loop BB0_11 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b64 v[57:58], v[53:54], off
	s_or_b32 s36, s36, exec_lo
	s_mov_b32 s37, exec_lo
                                        ; implicit-def: $vgpr40
	s_waitcnt vmcnt(0)
	v_cmpx_eq_u64_e64 v[57:58], v[6:7]
; %bb.15:                               ;   in Loop: Header=BB0_14 Depth=2
	v_add_nc_u32_e32 v40, 1, v56
	v_add_co_u32 v53, s3, v53, 8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e64 v54, s3, 0, v54, s3
	v_cmp_ge_i32_e32 vcc_lo, v40, v9
	v_mov_b32_e32 v56, v95
	s_and_not1_b32 s3, s36, exec_lo
	s_and_b32 s36, vcc_lo, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s36, s3, s36
; %bb.16:                               ; %Flow1013
                                        ;   in Loop: Header=BB0_14 Depth=2
	s_or_b32 exec_lo, exec_lo, s37
	v_mov_b32_e32 v55, v56
	v_mov_b32_e32 v56, v40
	s_and_b32 s3, exec_lo, s36
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s16, s3, s16
	s_and_not1_b32 exec_lo, exec_lo, s16
	s_cbranch_execnz .LBB0_14
; %bb.17:                               ;   in Loop: Header=BB0_11 Depth=1
	s_or_b32 exec_lo, exec_lo, s16
	s_mov_b32 s3, 0
	s_mov_b32 s56, exec_lo
	v_cmpx_ne_u32_e64 v55, v9
	s_cbranch_execz .LBB0_29
; %bb.18:                               ;   in Loop: Header=BB0_11 Depth=1
	v_ashrrev_i32_e32 v56, 31, v55
	v_mov_b32_e32 v69, 0
	v_mov_b32_e32 v70, 0x3ff00000
	s_mov_b64 s[36:37], 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[53:54], 3, v[55:56]
	v_add_co_u32 v53, vcc_lo, s50, v53
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_4) | instid1(VALU_DEP_2)
	v_add_co_ci_u32_e32 v54, vcc_lo, s51, v54, vcc_lo
	global_load_b64 v[57:58], v[53:54], off
	s_waitcnt vmcnt(0)
	v_lshlrev_b32_e32 v53, 1, v57
	v_ashrrev_i32_e32 v58, 31, v57
	v_ashrrev_i32_e32 v54, 31, v53
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[53:54], 3, v[53:54]
	v_add_co_u32 v53, vcc_lo, s6, v53
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_4) | instid1(VALU_DEP_2)
	v_add_co_ci_u32_e32 v54, vcc_lo, s7, v54, vcc_lo
	global_load_b128 v[53:56], v[53:54], off
	s_waitcnt vmcnt(0)
	v_add_f64 v[2:3], v[55:56], -v[2:3]
	v_add_f64 v[53:54], v[53:54], -v[0:1]
	v_mul_f64 v[0:1], v[2:3], v[2:3]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f64 v[59:60], v[53:54], v[53:54], v[0:1]
	v_lshlrev_b64 v[0:1], 3, v[57:58]
	v_add_co_u32 v57, vcc_lo, s8, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v58, vcc_lo, s9, v1, vcc_lo
	s_clause 0x1
	global_load_b64 v[55:56], v[17:18], off
	global_load_b64 v[57:58], v[57:58], off
	v_div_scale_f64 v[61:62], null, v[59:60], v[59:60], 1.0
	v_div_scale_f64 v[67:68], vcc_lo, 1.0, v[59:60], 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f64_e32 v[63:64], v[61:62]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[65:66], -v[61:62], v[63:64], 1.0
	v_fma_f64 v[63:64], v[63:64], v[65:66], v[63:64]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[65:66], -v[61:62], v[63:64], 1.0
	v_fma_f64 v[63:64], v[63:64], v[65:66], v[63:64]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[65:66], v[67:68], v[63:64]
	v_fma_f64 v[61:62], -v[61:62], v[65:66], v[67:68]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_div_fmas_f64 v[61:62], v[61:62], v[63:64], v[65:66]
	s_waitcnt vmcnt(0)
	v_add_f64 v[65:66], v[57:58], -v[55:56]
	v_div_fixup_f64 v[59:60], v[61:62], v[59:60], 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_nle_f64_e32 vcc_lo, 0, v[65:66]
	v_mul_f64 v[2:3], v[2:3], v[59:60]
	v_mul_f64 v[53:54], v[53:54], v[59:60]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f64 v[57:58], v[65:66], v[2:3]
	v_mul_f64 v[63:64], v[65:66], v[53:54]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_mul_f64 v[59:60], v[41:42], v[57:58]
	v_mul_f64 v[67:68], v[43:44], v[57:58]
	v_mul_f64 v[57:58], v[45:46], v[57:58]
	v_fma_f64 v[61:62], v[47:48], v[63:64], v[59:60]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[59:60], v[49:50], v[63:64], v[67:68]
	v_fma_f64 v[57:58], v[51:52], v[63:64], v[57:58]
	v_max_f64 v[63:64], v[65:66], 0
	v_dual_cndmask_b32 v66, 0, v66 :: v_dual_mov_b32 v67, 0
	v_dual_mov_b32 v68, 0x408f4000 :: v_dual_cndmask_b32 v65, 0, v65
.LBB0_19:                               ;   Parent Loop BB0_11 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_cmp_eq_u32 s36, 1
	s_mov_b32 s16, s4
	s_cselect_b32 vcc_lo, -1, 0
	s_cmp_eq_u32 s36, 2
	v_dual_cndmask_b32 v40, v62, v60 :: v_dual_cndmask_b32 v71, v61, v59
	s_cselect_b32 vcc_lo, -1, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_max_f64 v[67:68], v[67:68], v[67:68]
	s_add_u32 s36, s36, 1
	v_dual_cndmask_b32 v72, v40, v58 :: v_dual_cndmask_b32 v71, v71, v57
	s_addc_u32 s37, s37, 0
	s_cmp_lg_u32 s36, 3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_scale_f64 v[73:74], null, v[71:72], v[71:72], v[65:66]
	v_div_scale_f64 v[75:76], null, v[71:72], v[71:72], v[63:64]
	v_rcp_f64_e32 v[77:78], v[73:74]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_rcp_f64_e32 v[79:80], v[75:76]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[81:82], -v[73:74], v[77:78], 1.0
	v_fma_f64 v[83:84], -v[75:76], v[79:80], 1.0
	v_fma_f64 v[77:78], v[77:78], v[81:82], v[77:78]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_fma_f64 v[79:80], v[79:80], v[83:84], v[79:80]
	v_div_scale_f64 v[83:84], vcc_lo, v[65:66], v[71:72], v[65:66]
	v_fma_f64 v[81:82], -v[73:74], v[77:78], 1.0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[85:86], -v[75:76], v[79:80], 1.0
	v_fma_f64 v[77:78], v[77:78], v[81:82], v[77:78]
	v_div_scale_f64 v[81:82], s3, v[63:64], v[71:72], v[63:64]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[79:80], v[79:80], v[85:86], v[79:80]
	v_mul_f64 v[85:86], v[83:84], v[77:78]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f64 v[87:88], v[81:82], v[79:80]
	v_fma_f64 v[73:74], -v[73:74], v[85:86], v[83:84]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[75:76], -v[75:76], v[87:88], v[81:82]
	v_div_fmas_f64 v[73:74], v[73:74], v[77:78], v[85:86]
	s_mov_b32 vcc_lo, s3
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_div_fmas_f64 v[75:76], v[75:76], v[79:80], v[87:88]
	v_cmp_gt_f64_e32 vcc_lo, s[4:5], v[71:72]
	v_div_fixup_f64 v[73:74], v[73:74], v[71:72], v[65:66]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_div_fixup_f64 v[75:76], v[75:76], v[71:72], v[63:64]
	v_dual_cndmask_b32 v40, v69, v73 :: v_dual_cndmask_b32 v69, v70, v74
	v_cmp_lt_f64_e32 vcc_lo, s[16:17], v[71:72]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v70, v69, v76 :: v_dual_cndmask_b32 v69, v40, v75
	v_max_f64 v[71:72], v[69:70], v[69:70]
	s_delay_alu instid0(VALU_DEP_1)
	v_min_f64 v[67:68], v[71:72], v[67:68]
	s_cbranch_scc1 .LBB0_19
; %bb.20:                               ; %_Z14limit_gradientPdddd.exit733
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_3)
	v_mul_f64 v[63:64], v[67:68], s[38:39]
	s_mov_b64 s[36:37], 0
	v_mov_b32_e32 v65, 0
	v_mov_b32_e32 v66, 0x3ff00000
	v_min_f64 v[63:64], v[63:64], 1.0
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f64 v[55:56], v[61:62], v[63:64], v[55:56]
	global_store_b64 v[19:20], v[55:56], off
	global_load_b64 v[55:56], v[17:18], off
	s_waitcnt vmcnt(0)
	v_fma_f64 v[55:56], v[59:60], v[63:64], v[55:56]
	global_store_b64 v[19:20], v[55:56], off offset:8
	global_load_b64 v[55:56], v[17:18], off
	s_waitcnt vmcnt(0)
	v_fma_f64 v[55:56], v[57:58], v[63:64], v[55:56]
	v_mov_b32_e32 v63, 0
	v_add_co_u32 v57, vcc_lo, s12, v0
	v_add_co_ci_u32_e32 v58, vcc_lo, s13, v1, vcc_lo
	global_store_b64 v[19:20], v[55:56], off offset:16
	s_clause 0x1
	global_load_b64 v[55:56], v[57:58], off
	global_load_b64 v[57:58], v[12:13], off
	s_waitcnt vmcnt(0)
	v_add_f64 v[61:62], v[55:56], -v[57:58]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_4)
	v_mul_f64 v[53:54], v[53:54], v[61:62]
	v_mul_f64 v[59:60], v[2:3], v[61:62]
	v_cmp_nle_f64_e32 vcc_lo, 0, v[61:62]
	v_mov_b32_e32 v64, 0x408f4000
	v_mul_f64 v[47:48], v[47:48], v[53:54]
	v_mul_f64 v[49:50], v[49:50], v[53:54]
	v_mul_f64 v[51:52], v[51:52], v[53:54]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[57:58], v[41:42], v[59:60], v[47:48]
	v_fma_f64 v[55:56], v[43:44], v[59:60], v[49:50]
	s_delay_alu instid0(VALU_DEP_3)
	v_fma_f64 v[53:54], v[45:46], v[59:60], v[51:52]
	v_max_f64 v[59:60], v[61:62], 0
	v_dual_cndmask_b32 v61, 0, v61 :: v_dual_cndmask_b32 v62, 0, v62
.LBB0_21:                               ;   Parent Loop BB0_11 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_cmp_eq_u32 s36, 1
	s_mov_b32 s16, s4
	s_cselect_b32 vcc_lo, -1, 0
	s_cmp_eq_u32 s36, 2
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_dual_cndmask_b32 v40, v58, v56 :: v_dual_cndmask_b32 v67, v57, v55
	s_cselect_b32 vcc_lo, -1, 0
	v_max_f64 v[63:64], v[63:64], v[63:64]
	s_add_u32 s36, s36, 1
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v68, v40, v54 :: v_dual_cndmask_b32 v67, v67, v53
	s_addc_u32 s37, s37, 0
	s_cmp_lg_u32 s36, 3
	v_div_scale_f64 v[69:70], null, v[67:68], v[67:68], v[61:62]
	v_div_scale_f64 v[71:72], null, v[67:68], v[67:68], v[59:60]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_rcp_f64_e32 v[73:74], v[69:70]
	v_rcp_f64_e32 v[75:76], v[71:72]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[77:78], -v[69:70], v[73:74], 1.0
	v_fma_f64 v[79:80], -v[71:72], v[75:76], 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[73:74], v[73:74], v[77:78], v[73:74]
	v_fma_f64 v[75:76], v[75:76], v[79:80], v[75:76]
	v_div_scale_f64 v[79:80], vcc_lo, v[61:62], v[67:68], v[61:62]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[77:78], -v[69:70], v[73:74], 1.0
	v_fma_f64 v[81:82], -v[71:72], v[75:76], 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_fma_f64 v[73:74], v[73:74], v[77:78], v[73:74]
	v_div_scale_f64 v[77:78], s3, v[59:60], v[67:68], v[59:60]
	v_fma_f64 v[75:76], v[75:76], v[81:82], v[75:76]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f64 v[81:82], v[79:80], v[73:74]
	v_mul_f64 v[83:84], v[77:78], v[75:76]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[69:70], -v[69:70], v[81:82], v[79:80]
	v_fma_f64 v[71:72], -v[71:72], v[83:84], v[77:78]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_fmas_f64 v[69:70], v[69:70], v[73:74], v[81:82]
	s_mov_b32 vcc_lo, s3
	v_div_fmas_f64 v[71:72], v[71:72], v[75:76], v[83:84]
	v_cmp_gt_f64_e32 vcc_lo, s[4:5], v[67:68]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_div_fixup_f64 v[69:70], v[69:70], v[67:68], v[61:62]
	v_div_fixup_f64 v[71:72], v[71:72], v[67:68], v[59:60]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_dual_cndmask_b32 v40, v65, v69 :: v_dual_cndmask_b32 v65, v66, v70
	v_cmp_lt_f64_e32 vcc_lo, s[16:17], v[67:68]
	v_dual_cndmask_b32 v66, v65, v72 :: v_dual_cndmask_b32 v65, v40, v71
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_max_f64 v[67:68], v[65:66], v[65:66]
	v_min_f64 v[63:64], v[67:68], v[63:64]
	s_cbranch_scc1 .LBB0_21
; %bb.22:                               ; %_Z14limit_gradientPdddd.exit740
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mul_f64 v[59:60], v[63:64], s[38:39]
	s_mov_b64 s[36:37], 0
	v_min_f64 v[59:60], v[59:60], 1.0
	s_delay_alu instid0(VALU_DEP_1)
	v_mul_f64 v[57:58], v[57:58], v[59:60]
	v_mul_f64 v[55:56], v[55:56], v[59:60]
	v_mul_f64 v[53:54], v[53:54], v[59:60]
	v_dual_mov_b32 v60, v32 :: v_dual_mov_b32 v59, v31
.LBB0_23:                               ;   Parent Loop BB0_11 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b64 v[61:62], v[12:13], off
	s_cmp_eq_u32 s36, 1
	s_cselect_b32 vcc_lo, -1, 0
	s_cmp_eq_u32 s36, 2
	v_dual_cndmask_b32 v40, v58, v56 :: v_dual_cndmask_b32 v63, v57, v55
	s_cselect_b32 vcc_lo, -1, 0
	s_add_u32 s36, s36, 1
	s_addc_u32 s37, s37, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v64, v40, v54 :: v_dual_cndmask_b32 v63, v63, v53
	s_cmp_lg_u32 s36, 3
	s_waitcnt vmcnt(0)
	v_add_f64 v[61:62], v[63:64], v[61:62]
	global_store_b64 v[59:60], v[61:62], off
	v_add_co_u32 v59, vcc_lo, v59, 8
	v_add_co_ci_u32_e32 v60, vcc_lo, 0, v60, vcc_lo
	s_cbranch_scc1 .LBB0_23
; %bb.24:                               ;   in Loop: Header=BB0_11 Depth=1
	v_add_co_u32 v0, vcc_lo, s14, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s15, v1, vcc_lo
	s_mov_b64 s[36:37], 0
	s_clause 0x1
	global_load_b64 v[0:1], v[0:1], off
	global_load_b64 v[53:54], v[10:11], off
	s_waitcnt vmcnt(0)
	v_add_f64 v[53:54], v[0:1], -v[53:54]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_f64 v[0:1], v[2:3], v[53:54]
	v_cmp_nle_f64_e32 vcc_lo, 0, v[53:54]
	v_fma_f64 v[40:41], v[41:42], v[0:1], v[47:48]
	v_fma_f64 v[2:3], v[43:44], v[0:1], v[49:50]
	v_fma_f64 v[0:1], v[45:46], v[0:1], v[51:52]
	v_max_f64 v[42:43], v[53:54], 0
	v_dual_cndmask_b32 v45, 0, v54 :: v_dual_mov_b32 v48, 0
	v_dual_mov_b32 v49, 0x3ff00000 :: v_dual_mov_b32 v46, 0
	v_dual_mov_b32 v47, 0x408f4000 :: v_dual_cndmask_b32 v44, 0, v53
.LBB0_25:                               ;   Parent Loop BB0_11 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_cmp_eq_u32 s36, 1
	s_mov_b32 s16, s4
	s_cselect_b32 vcc_lo, -1, 0
	s_cmp_eq_u32 s36, 2
	v_cndmask_b32_e32 v50, v41, v3, vcc_lo
	v_cndmask_b32_e32 v52, v40, v2, vcc_lo
	s_cselect_b32 vcc_lo, -1, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_max_f64 v[46:47], v[46:47], v[46:47]
	s_add_u32 s36, s36, 1
	v_dual_cndmask_b32 v51, v50, v1 :: v_dual_cndmask_b32 v50, v52, v0
	s_addc_u32 s37, s37, 0
	s_cmp_lg_u32 s36, 3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_scale_f64 v[52:53], null, v[50:51], v[50:51], v[44:45]
	v_div_scale_f64 v[54:55], null, v[50:51], v[50:51], v[42:43]
	v_rcp_f64_e32 v[56:57], v[52:53]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_rcp_f64_e32 v[58:59], v[54:55]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[60:61], -v[52:53], v[56:57], 1.0
	v_fma_f64 v[62:63], -v[54:55], v[58:59], 1.0
	v_fma_f64 v[56:57], v[56:57], v[60:61], v[56:57]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_fma_f64 v[58:59], v[58:59], v[62:63], v[58:59]
	v_div_scale_f64 v[62:63], vcc_lo, v[44:45], v[50:51], v[44:45]
	v_fma_f64 v[60:61], -v[52:53], v[56:57], 1.0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[64:65], -v[54:55], v[58:59], 1.0
	v_fma_f64 v[56:57], v[56:57], v[60:61], v[56:57]
	v_div_scale_f64 v[60:61], s3, v[42:43], v[50:51], v[42:43]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[58:59], v[58:59], v[64:65], v[58:59]
	v_mul_f64 v[64:65], v[62:63], v[56:57]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f64 v[66:67], v[60:61], v[58:59]
	v_fma_f64 v[52:53], -v[52:53], v[64:65], v[62:63]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[54:55], -v[54:55], v[66:67], v[60:61]
	v_div_fmas_f64 v[52:53], v[52:53], v[56:57], v[64:65]
	s_mov_b32 vcc_lo, s3
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_div_fmas_f64 v[54:55], v[54:55], v[58:59], v[66:67]
	v_cmp_gt_f64_e32 vcc_lo, s[4:5], v[50:51]
	v_div_fixup_f64 v[52:53], v[52:53], v[50:51], v[44:45]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e32 v49, v49, v53, vcc_lo
	v_div_fixup_f64 v[54:55], v[54:55], v[50:51], v[42:43]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v48, v48, v52, vcc_lo
	v_cmp_lt_f64_e32 vcc_lo, s[16:17], v[50:51]
	v_dual_cndmask_b32 v49, v49, v55 :: v_dual_cndmask_b32 v48, v48, v54
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_max_f64 v[50:51], v[48:49], v[48:49]
	v_min_f64 v[46:47], v[50:51], v[46:47]
	s_cbranch_scc1 .LBB0_25
; %bb.26:                               ; %_Z14limit_gradientPdddd.exit747
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mul_f64 v[42:43], v[46:47], s[38:39]
	s_mov_b64 s[36:37], 0
	v_min_f64 v[42:43], v[42:43], 1.0
	s_delay_alu instid0(VALU_DEP_1)
	v_mul_f64 v[40:41], v[40:41], v[42:43]
	v_mul_f64 v[2:3], v[2:3], v[42:43]
	v_mul_f64 v[0:1], v[0:1], v[42:43]
	v_dual_mov_b32 v43, v34 :: v_dual_mov_b32 v42, v33
.LBB0_27:                               ;   Parent Loop BB0_11 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b64 v[44:45], v[10:11], off
	s_cmp_eq_u32 s36, 1
	s_cselect_b32 vcc_lo, -1, 0
	s_cmp_eq_u32 s36, 2
	v_cndmask_b32_e32 v46, v41, v3, vcc_lo
	v_cndmask_b32_e32 v48, v40, v2, vcc_lo
	s_cselect_b32 vcc_lo, -1, 0
	s_add_u32 s36, s36, 1
	s_addc_u32 s37, s37, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v47, v46, v1 :: v_dual_cndmask_b32 v46, v48, v0
	s_cmp_eq_u32 s36, 3
	s_waitcnt vmcnt(0)
	v_add_f64 v[44:45], v[46:47], v[44:45]
	global_store_b64 v[42:43], v[44:45], off
	v_add_co_u32 v42, vcc_lo, v42, 8
	v_add_co_ci_u32_e32 v43, vcc_lo, 0, v43, vcc_lo
	s_cbranch_scc0 .LBB0_27
; %bb.28:                               ; %Flow1012
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_mov_b32 s3, exec_lo
.LBB0_29:                               ; %Flow1021
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_or_b32 exec_lo, exec_lo, s56
	s_mov_b32 s56, -1
	s_and_b32 s57, s3, exec_lo
                                        ; implicit-def: $vgpr41_vgpr42
                                        ; implicit-def: $vgpr47_vgpr48
                                        ; implicit-def: $vgpr43_vgpr44
                                        ; implicit-def: $vgpr49_vgpr50
                                        ; implicit-def: $vgpr45_vgpr46
                                        ; implicit-def: $vgpr51_vgpr52
.LBB0_30:                               ; %Flow1020
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_or_saveexec_b32 s58, s58
                                        ; implicit-def: $vgpr58_vgpr59
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 exec_lo, exec_lo, s58
	s_cbranch_execz .LBB0_46
; %bb.31:                               ;   in Loop: Header=BB0_11 Depth=1
	s_clause 0x1
	global_load_b128 v[0:3], v[27:28], off
	global_load_b32 v53, v[27:28], off offset:16
	s_mov_b32 s59, 0
	s_mov_b32 s3, exec_lo
	s_waitcnt vmcnt(1)
	v_ashrrev_i32_e32 v3, 31, v2
	s_waitcnt vmcnt(0)
	v_ashrrev_i32_e32 v54, 31, v53
	v_ashrrev_i32_e32 v1, 31, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[55:56], 3, v[2:3]
	v_lshlrev_b64 v[57:58], 3, v[53:54]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[63:64], 3, v[0:1]
	v_add_co_u32 v59, vcc_lo, s10, v55
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_ci_u32_e32 v60, vcc_lo, s11, v56, vcc_lo
	v_add_co_u32 v61, vcc_lo, s10, v57
	v_add_co_ci_u32_e32 v62, vcc_lo, s11, v58, vcc_lo
	v_add_co_u32 v54, vcc_lo, s8, v55
	v_add_co_ci_u32_e32 v55, vcc_lo, s9, v56, vcc_lo
	v_add_co_u32 v56, vcc_lo, s8, v57
	v_add_co_ci_u32_e32 v57, vcc_lo, s9, v58, vcc_lo
	s_clause 0x1
	global_load_b64 v[61:62], v[61:62], off
	global_load_b64 v[58:59], v[59:60], off
	s_clause 0x1
	global_load_b64 v[54:55], v[54:55], off
	global_load_b64 v[56:57], v[56:57], off
	v_add_co_u32 v65, vcc_lo, s10, v63
	v_add_co_ci_u32_e32 v66, vcc_lo, s11, v64, vcc_lo
	v_add_co_u32 v63, vcc_lo, s8, v63
	v_add_co_ci_u32_e32 v64, vcc_lo, s9, v64, vcc_lo
	global_load_b64 v[69:70], v[65:66], off
	s_clause 0x1
	global_load_b64 v[63:64], v[63:64], off
	global_load_b64 v[65:66], v[17:18], off
	global_load_b64 v[67:68], v[29:30], off
	s_waitcnt vmcnt(7)
	v_max_f64 v[60:61], v[61:62], v[61:62]
	s_waitcnt vmcnt(6)
	v_max_f64 v[58:59], v[58:59], v[58:59]
	s_waitcnt vmcnt(5)
	v_max_f64 v[54:55], v[54:55], v[54:55]
	s_waitcnt vmcnt(4)
	v_max_f64 v[56:57], v[56:57], v[56:57]
	s_waitcnt vmcnt(3)
	v_max_f64 v[69:70], v[69:70], v[69:70]
	s_waitcnt vmcnt(2)
	v_max_f64 v[62:63], v[63:64], v[63:64]
	s_waitcnt vmcnt(1)
	v_max_f64 v[71:72], v[65:66], v[65:66]
	s_waitcnt vmcnt(0)
	v_max_f64 v[73:74], v[67:68], v[67:68]
	v_max_f64 v[54:55], v[54:55], v[58:59]
	v_max_f64 v[56:57], v[56:57], v[60:61]
	v_max_f64 v[58:59], v[58:59], v[60:61]
	v_max_f64 v[60:61], v[62:63], v[69:70]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_min_f64 v[54:55], v[54:55], v[56:57]
	v_max_f64 v[56:57], v[71:72], v[73:74]
	v_max_f64 v[58:59], v[69:70], v[58:59]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_min_f64 v[54:55], v[60:61], v[54:55]
	v_max_f64 v[58:59], v[73:74], v[58:59]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_min_f64 v[54:55], v[56:57], v[54:55]
	v_cmp_lt_f64_e32 vcc_lo, v[54:55], v[58:59]
	v_cndmask_b32_e32 v57, v0, v6, vcc_lo
	v_cndmask_b32_e32 v61, v53, v6, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_dual_cndmask_b32 v69, v2, v6 :: v_dual_lshlrev_b32 v2, 1, v57
	v_lshlrev_b32_e32 v0, 1, v61
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b32_e32 v53, 1, v69
	v_ashrrev_i32_e32 v3, 31, v2
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_ashrrev_i32_e32 v1, 31, v0
	v_ashrrev_i32_e32 v54, 31, v53
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[2:3], 3, v[2:3]
	v_lshlrev_b64 v[0:1], 3, v[0:1]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[53:54], 3, v[53:54]
	v_add_co_u32 v0, vcc_lo, s6, v0
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v1, vcc_lo, s7, v1, vcc_lo
	v_add_co_u32 v55, vcc_lo, s6, v2
	v_add_co_ci_u32_e32 v56, vcc_lo, s7, v3, vcc_lo
	v_add_co_u32 v53, vcc_lo, s6, v53
	v_add_co_ci_u32_e32 v54, vcc_lo, s7, v54, vcc_lo
	s_clause 0x2
	global_load_b128 v[0:3], v[0:1], off
	global_load_b128 v[70:73], v[55:56], off
	global_load_b128 v[74:77], v[53:54], off
	s_waitcnt vmcnt(1)
	v_add_f64 v[53:54], v[0:1], -v[70:71]
	s_waitcnt vmcnt(0)
	v_add_f64 v[55:56], v[76:77], -v[72:73]
	v_add_f64 v[0:1], v[74:75], -v[70:71]
	v_add_f64 v[2:3], v[2:3], -v[72:73]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[58:59], v[55:56], v[53:54]
	v_fma_f64 v[63:64], v[0:1], v[2:3], -v[58:59]
                                        ; implicit-def: $vgpr58_vgpr59
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_ge_f64_e32 0, v[63:64]
	s_xor_b32 s3, exec_lo, s3
	s_cbranch_execz .LBB0_33
; %bb.32:                               ;   in Loop: Header=BB0_11 Depth=1
	global_store_b64 v[19:20], v[65:66], off
	global_load_b64 v[58:59], v[17:18], off
	s_mov_b32 s59, exec_lo
                                        ; implicit-def: $vgpr67_vgpr68
                                        ; implicit-def: $vgpr65_vgpr66
                                        ; implicit-def: $vgpr61
                                        ; implicit-def: $vgpr69
                                        ; implicit-def: $vgpr57
                                        ; implicit-def: $vgpr0_vgpr1
                                        ; implicit-def: $vgpr53_vgpr54
                                        ; implicit-def: $vgpr55_vgpr56
                                        ; implicit-def: $vgpr2_vgpr3
                                        ; implicit-def: $vgpr63_vgpr64
                                        ; implicit-def: $vgpr41_vgpr42
                                        ; implicit-def: $vgpr47_vgpr48
                                        ; implicit-def: $vgpr43_vgpr44
                                        ; implicit-def: $vgpr49_vgpr50
                                        ; implicit-def: $vgpr45_vgpr46
                                        ; implicit-def: $vgpr51_vgpr52
	s_waitcnt vmcnt(0)
	global_store_b64 v[19:20], v[58:59], off offset:8
.LBB0_33:                               ; %Flow1023
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_or_saveexec_b32 s60, s3
	s_mov_b32 s3, s57
	s_xor_b32 exec_lo, exec_lo, s60
	s_cbranch_execz .LBB0_45
; %bb.34:                               ;   in Loop: Header=BB0_11 Depth=1
	v_ashrrev_i32_e32 v58, 31, v57
	v_ashrrev_i32_e32 v70, 31, v69
	v_ashrrev_i32_e32 v62, 31, v61
	v_div_scale_f64 v[81:82], null, v[63:64], v[63:64], 1.0
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_lshlrev_b64 v[57:58], 3, v[57:58]
	v_lshlrev_b64 v[59:60], 3, v[69:70]
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[61:62], 3, v[61:62]
	s_mov_b64 s[36:37], 0
	v_add_co_u32 v69, vcc_lo, s8, v57
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_ci_u32_e32 v70, vcc_lo, s9, v58, vcc_lo
	v_add_co_u32 v71, vcc_lo, s8, v59
	v_add_co_ci_u32_e32 v72, vcc_lo, s9, v60, vcc_lo
	v_add_co_u32 v73, vcc_lo, s8, v61
	s_clause 0x1
	global_load_b64 v[69:70], v[69:70], off
	global_load_b64 v[71:72], v[71:72], off
	v_add_co_ci_u32_e32 v74, vcc_lo, s9, v62, vcc_lo
	global_load_b64 v[73:74], v[73:74], off
	v_rcp_f64_e32 v[75:76], v[81:82]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[77:78], -v[81:82], v[75:76], 1.0
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f64 v[83:84], v[75:76], v[77:78], v[75:76]
	v_add_co_u32 v75, vcc_lo, s10, v57
	v_add_co_ci_u32_e32 v76, vcc_lo, s11, v58, vcc_lo
	v_add_co_u32 v77, vcc_lo, s10, v59
	v_add_co_ci_u32_e32 v78, vcc_lo, s11, v60, vcc_lo
	v_add_co_u32 v79, vcc_lo, s10, v61
	v_add_co_ci_u32_e32 v80, vcc_lo, s11, v62, vcc_lo
	s_clause 0x2
	global_load_b64 v[75:76], v[75:76], off
	global_load_b64 v[77:78], v[77:78], off
	global_load_b64 v[79:80], v[79:80], off
	v_div_scale_f64 v[87:88], vcc_lo, 1.0, v[63:64], 1.0
	v_fma_f64 v[85:86], -v[81:82], v[83:84], 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[83:84], v[83:84], v[85:86], v[83:84]
	v_mul_f64 v[85:86], v[87:88], v[83:84]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[81:82], -v[81:82], v[85:86], v[87:88]
	v_div_fmas_f64 v[81:82], v[81:82], v[83:84], v[85:86]
	s_waitcnt vmcnt(5)
	v_add_f64 v[85:86], v[69:70], -v[65:66]
	s_waitcnt vmcnt(4)
	v_add_f64 v[87:88], v[71:72], -v[69:70]
	s_waitcnt vmcnt(3)
	v_add_f64 v[83:84], v[73:74], -v[69:70]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_div_fixup_f64 v[63:64], v[81:82], v[63:64], 1.0
	v_mul_f64 v[89:90], v[53:54], v[87:88]
	v_add_f64 v[91:92], v[85:86], v[87:88]
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_mul_f64 v[81:82], v[55:56], v[83:84]
	v_add_f64 v[93:94], v[85:86], v[83:84]
	v_fma_f64 v[83:84], v[0:1], v[83:84], -v[89:90]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[81:82], v[2:3], v[87:88], -v[81:82]
	v_max_f64 v[87:88], v[91:92], v[93:94]
	v_min_f64 v[89:90], v[91:92], v[93:94]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_mul_f64 v[83:84], v[63:64], v[83:84]
	v_mul_f64 v[81:82], v[63:64], v[81:82]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_max_f64 v[87:88], v[85:86], v[87:88]
	v_min_f64 v[98:99], v[85:86], v[89:90]
	s_delay_alu instid0(VALU_DEP_4)
	v_mul_f64 v[91:92], v[41:42], v[83:84]
	v_mul_f64 v[93:94], v[43:44], v[83:84]
	v_mul_f64 v[96:97], v[45:46], v[83:84]
	v_max_f64 v[89:90], v[87:88], 0
	v_mov_b32_e32 v87, 0
	v_mov_b32_e32 v88, 0x408f4000
	v_fma_f64 v[85:86], v[47:48], v[81:82], v[91:92]
	v_fma_f64 v[83:84], v[49:50], v[81:82], v[93:94]
	v_fma_f64 v[81:82], v[51:52], v[81:82], v[96:97]
	v_min_f64 v[91:92], v[98:99], 0
	v_mov_b32_e32 v93, 0
	v_mov_b32_e32 v94, 0x3ff00000
.LBB0_35:                               ;   Parent Loop BB0_11 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_cmp_eq_u32 s36, 1
	s_mov_b32 s16, s4
	s_cselect_b32 vcc_lo, -1, 0
	s_cmp_eq_u32 s36, 2
	v_cndmask_b32_e32 v40, v86, v84, vcc_lo
	v_cndmask_b32_e32 v96, v85, v83, vcc_lo
	s_cselect_b32 vcc_lo, -1, 0
	s_delay_alu instid0(VALU_DEP_3)
	v_max_f64 v[87:88], v[87:88], v[87:88]
	s_add_u32 s36, s36, 1
	v_cndmask_b32_e32 v97, v40, v82, vcc_lo
	v_cndmask_b32_e32 v96, v96, v81, vcc_lo
	s_addc_u32 s37, s37, 0
	s_cmp_lg_u32 s36, 3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_scale_f64 v[98:99], null, v[96:97], v[96:97], v[91:92]
	v_div_scale_f64 v[100:101], null, v[96:97], v[96:97], v[89:90]
	v_rcp_f64_e32 v[102:103], v[98:99]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_rcp_f64_e32 v[104:105], v[100:101]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[106:107], -v[98:99], v[102:103], 1.0
	v_fma_f64 v[108:109], -v[100:101], v[104:105], 1.0
	v_fma_f64 v[102:103], v[102:103], v[106:107], v[102:103]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_fma_f64 v[104:105], v[104:105], v[108:109], v[104:105]
	v_div_scale_f64 v[108:109], vcc_lo, v[91:92], v[96:97], v[91:92]
	v_fma_f64 v[106:107], -v[98:99], v[102:103], 1.0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[110:111], -v[100:101], v[104:105], 1.0
	v_fma_f64 v[102:103], v[102:103], v[106:107], v[102:103]
	v_div_scale_f64 v[106:107], s3, v[89:90], v[96:97], v[89:90]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[104:105], v[104:105], v[110:111], v[104:105]
	v_mul_f64 v[110:111], v[108:109], v[102:103]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f64 v[112:113], v[106:107], v[104:105]
	v_fma_f64 v[98:99], -v[98:99], v[110:111], v[108:109]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[100:101], -v[100:101], v[112:113], v[106:107]
	v_div_fmas_f64 v[98:99], v[98:99], v[102:103], v[110:111]
	s_mov_b32 vcc_lo, s3
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_div_fmas_f64 v[100:101], v[100:101], v[104:105], v[112:113]
	v_cmp_gt_f64_e32 vcc_lo, s[4:5], v[96:97]
	v_div_fixup_f64 v[98:99], v[98:99], v[96:97], v[91:92]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_dual_cndmask_b32 v40, v93, v98 :: v_dual_cndmask_b32 v93, v94, v99
	v_div_fixup_f64 v[100:101], v[100:101], v[96:97], v[89:90]
	v_cmp_lt_f64_e32 vcc_lo, s[16:17], v[96:97]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v94, v93, v101 :: v_dual_cndmask_b32 v93, v40, v100
	v_max_f64 v[96:97], v[93:94], v[93:94]
	s_delay_alu instid0(VALU_DEP_1)
	v_min_f64 v[87:88], v[96:97], v[87:88]
	s_cbranch_scc1 .LBB0_35
; %bb.36:                               ; %_Z14limit_gradientPdddd.exit
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_waitcnt vmcnt(1)
	v_add_f64 v[71:72], v[71:72], -v[77:78]
	s_waitcnt vmcnt(0)
	v_add_f64 v[73:74], v[73:74], -v[79:80]
	v_add_f64 v[69:70], v[69:70], -v[75:76]
	v_add_f64 v[67:68], v[65:66], -v[67:68]
	s_mov_b64 s[36:37], 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_min_f64 v[71:72], v[71:72], v[73:74]
	v_min_f64 v[69:70], v[69:70], v[71:72]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_min_f64 v[67:68], v[69:70], v[67:68]
	v_cmp_lt_f64_e32 vcc_lo, 0, v[67:68]
	v_cndmask_b32_e64 v40, 0, 0x3ff00000, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[67:68], v[21:22], v[39:40], s[40:41]
	v_mul_f64 v[67:68], v[67:68], v[87:88]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_min_f64 v[67:68], v[67:68], 1.0
	v_fma_f64 v[65:66], v[85:86], v[67:68], v[65:66]
	global_store_b64 v[19:20], v[65:66], off
	global_load_b64 v[65:66], v[17:18], off
	s_waitcnt vmcnt(0)
	v_fma_f64 v[65:66], v[83:84], v[67:68], v[65:66]
	global_store_b64 v[19:20], v[65:66], off offset:8
	global_load_b64 v[65:66], v[17:18], off
	s_waitcnt vmcnt(0)
	v_fma_f64 v[65:66], v[81:82], v[67:68], v[65:66]
	v_add_co_u32 v67, vcc_lo, s12, v57
	v_add_co_ci_u32_e32 v68, vcc_lo, s13, v58, vcc_lo
	v_add_co_u32 v69, vcc_lo, s12, v59
	v_add_co_ci_u32_e32 v70, vcc_lo, s13, v60, vcc_lo
	global_store_b64 v[19:20], v[65:66], off offset:16
	s_clause 0x1
	global_load_b64 v[65:66], v[67:68], off
	global_load_b64 v[67:68], v[69:70], off
	v_add_co_u32 v69, vcc_lo, s12, v61
	v_add_co_ci_u32_e32 v70, vcc_lo, s13, v62, vcc_lo
	s_clause 0x1
	global_load_b64 v[69:70], v[69:70], off
	global_load_b64 v[71:72], v[12:13], off
	s_waitcnt vmcnt(2)
	v_add_f64 v[67:68], v[67:68], -v[65:66]
	s_waitcnt vmcnt(1)
	v_add_f64 v[69:70], v[69:70], -v[65:66]
	s_waitcnt vmcnt(0)
	v_add_f64 v[65:66], v[65:66], -v[71:72]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_f64 v[71:72], v[53:54], v[67:68]
	v_mul_f64 v[73:74], v[55:56], v[69:70]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_add_f64 v[75:76], v[65:66], v[67:68]
	v_add_f64 v[77:78], v[65:66], v[69:70]
	v_fma_f64 v[69:70], v[0:1], v[69:70], -v[71:72]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[67:68], v[2:3], v[67:68], -v[73:74]
	v_max_f64 v[71:72], v[75:76], v[77:78]
	v_min_f64 v[73:74], v[75:76], v[77:78]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_mul_f64 v[69:70], v[63:64], v[69:70]
	v_mul_f64 v[75:76], v[63:64], v[67:68]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_max_f64 v[67:68], v[65:66], v[71:72]
	v_min_f64 v[73:74], v[65:66], v[73:74]
	s_delay_alu instid0(VALU_DEP_4)
	v_mul_f64 v[77:78], v[41:42], v[69:70]
	v_mul_f64 v[79:80], v[43:44], v[69:70]
	v_mul_f64 v[81:82], v[45:46], v[69:70]
	v_max_f64 v[71:72], v[67:68], 0
	v_min_f64 v[73:74], v[73:74], 0
	v_fma_f64 v[69:70], v[47:48], v[75:76], v[77:78]
	v_fma_f64 v[67:68], v[49:50], v[75:76], v[79:80]
	v_fma_f64 v[65:66], v[51:52], v[75:76], v[81:82]
	v_mov_b32_e32 v77, 0
	v_dual_mov_b32 v78, 0x3ff00000 :: v_dual_mov_b32 v75, 0
	v_mov_b32_e32 v76, 0x408f4000
.LBB0_37:                               ;   Parent Loop BB0_11 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_cmp_eq_u32 s36, 1
	s_mov_b32 s16, s4
	s_cselect_b32 vcc_lo, -1, 0
	s_cmp_eq_u32 s36, 2
	v_cndmask_b32_e32 v79, v70, v68, vcc_lo
	v_cndmask_b32_e32 v81, v69, v67, vcc_lo
	s_cselect_b32 vcc_lo, -1, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_max_f64 v[75:76], v[75:76], v[75:76]
	s_add_u32 s36, s36, 1
	v_dual_cndmask_b32 v80, v79, v66 :: v_dual_cndmask_b32 v79, v81, v65
	s_addc_u32 s37, s37, 0
	s_cmp_lg_u32 s36, 3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_scale_f64 v[81:82], null, v[79:80], v[79:80], v[73:74]
	v_div_scale_f64 v[83:84], null, v[79:80], v[79:80], v[71:72]
	v_rcp_f64_e32 v[85:86], v[81:82]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_rcp_f64_e32 v[87:88], v[83:84]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[89:90], -v[81:82], v[85:86], 1.0
	v_fma_f64 v[91:92], -v[83:84], v[87:88], 1.0
	v_fma_f64 v[85:86], v[85:86], v[89:90], v[85:86]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_fma_f64 v[87:88], v[87:88], v[91:92], v[87:88]
	v_div_scale_f64 v[91:92], vcc_lo, v[73:74], v[79:80], v[73:74]
	v_fma_f64 v[89:90], -v[81:82], v[85:86], 1.0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[93:94], -v[83:84], v[87:88], 1.0
	v_fma_f64 v[85:86], v[85:86], v[89:90], v[85:86]
	v_div_scale_f64 v[89:90], s3, v[71:72], v[79:80], v[71:72]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[87:88], v[87:88], v[93:94], v[87:88]
	v_mul_f64 v[93:94], v[91:92], v[85:86]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f64 v[96:97], v[89:90], v[87:88]
	v_fma_f64 v[81:82], -v[81:82], v[93:94], v[91:92]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[83:84], -v[83:84], v[96:97], v[89:90]
	v_div_fmas_f64 v[81:82], v[81:82], v[85:86], v[93:94]
	s_mov_b32 vcc_lo, s3
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_div_fmas_f64 v[83:84], v[83:84], v[87:88], v[96:97]
	v_cmp_gt_f64_e32 vcc_lo, s[4:5], v[79:80]
	v_div_fixup_f64 v[81:82], v[81:82], v[79:80], v[73:74]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e32 v78, v78, v82, vcc_lo
	v_div_fixup_f64 v[83:84], v[83:84], v[79:80], v[71:72]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v77, v77, v81, vcc_lo
	v_cmp_lt_f64_e32 vcc_lo, s[16:17], v[79:80]
	v_dual_cndmask_b32 v78, v78, v84 :: v_dual_cndmask_b32 v77, v77, v83
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_max_f64 v[79:80], v[77:78], v[77:78]
	v_min_f64 v[75:76], v[79:80], v[75:76]
	s_cbranch_scc1 .LBB0_37
; %bb.38:                               ; %_Z14limit_gradientPdddd.exit719
                                        ;   in Loop: Header=BB0_11 Depth=1
	v_fma_f64 v[71:72], v[23:24], v[39:40], s[44:45]
	s_mov_b64 s[36:37], 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[71:72], v[71:72], v[75:76]
	v_min_f64 v[71:72], v[71:72], 1.0
	s_delay_alu instid0(VALU_DEP_1)
	v_mul_f64 v[69:70], v[69:70], v[71:72]
	v_mul_f64 v[67:68], v[67:68], v[71:72]
	v_mul_f64 v[65:66], v[65:66], v[71:72]
	v_dual_mov_b32 v72, v32 :: v_dual_mov_b32 v71, v31
.LBB0_39:                               ;   Parent Loop BB0_11 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b64 v[73:74], v[12:13], off
	s_cmp_eq_u32 s36, 1
	s_cselect_b32 vcc_lo, -1, 0
	s_cmp_eq_u32 s36, 2
	v_cndmask_b32_e32 v75, v70, v68, vcc_lo
	v_cndmask_b32_e32 v77, v69, v67, vcc_lo
	s_cselect_b32 vcc_lo, -1, 0
	s_add_u32 s36, s36, 1
	s_addc_u32 s37, s37, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v76, v75, v66 :: v_dual_cndmask_b32 v75, v77, v65
	s_cmp_lg_u32 s36, 3
	s_waitcnt vmcnt(0)
	v_add_f64 v[73:74], v[75:76], v[73:74]
	global_store_b64 v[71:72], v[73:74], off
	v_add_co_u32 v71, vcc_lo, v71, 8
	v_add_co_ci_u32_e32 v72, vcc_lo, 0, v72, vcc_lo
	s_cbranch_scc1 .LBB0_39
; %bb.40:                               ;   in Loop: Header=BB0_11 Depth=1
	v_add_co_u32 v57, vcc_lo, s14, v57
	v_add_co_ci_u32_e32 v58, vcc_lo, s15, v58, vcc_lo
	v_add_co_u32 v59, vcc_lo, s14, v59
	v_add_co_ci_u32_e32 v60, vcc_lo, s15, v60, vcc_lo
	v_add_co_u32 v61, vcc_lo, s14, v61
	s_clause 0x1
	global_load_b64 v[57:58], v[57:58], off
	global_load_b64 v[59:60], v[59:60], off
	v_add_co_ci_u32_e32 v62, vcc_lo, s15, v62, vcc_lo
	s_clause 0x1
	global_load_b64 v[61:62], v[61:62], off
	global_load_b64 v[65:66], v[10:11], off
	s_mov_b64 s[36:37], 0
	s_waitcnt vmcnt(2)
	v_add_f64 v[59:60], v[59:60], -v[57:58]
	s_waitcnt vmcnt(1)
	v_add_f64 v[61:62], v[61:62], -v[57:58]
	s_waitcnt vmcnt(0)
	v_add_f64 v[57:58], v[57:58], -v[65:66]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_f64 v[53:54], v[53:54], v[59:60]
	v_mul_f64 v[55:56], v[55:56], v[61:62]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_add_f64 v[65:66], v[57:58], v[59:60]
	v_add_f64 v[67:68], v[57:58], v[61:62]
	v_fma_f64 v[0:1], v[0:1], v[61:62], -v[53:54]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[2:3], v[2:3], v[59:60], -v[55:56]
	v_max_f64 v[53:54], v[65:66], v[67:68]
	v_min_f64 v[55:56], v[65:66], v[67:68]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_mul_f64 v[0:1], v[63:64], v[0:1]
	v_mul_f64 v[59:60], v[63:64], v[2:3]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_max_f64 v[2:3], v[57:58], v[53:54]
	v_mul_f64 v[41:42], v[41:42], v[0:1]
	v_mul_f64 v[53:54], v[43:44], v[0:1]
	v_mul_f64 v[0:1], v[45:46], v[0:1]
	v_min_f64 v[45:46], v[57:58], v[55:56]
	v_max_f64 v[43:44], v[2:3], 0
	v_fma_f64 v[41:42], v[47:48], v[59:60], v[41:42]
	v_fma_f64 v[2:3], v[49:50], v[59:60], v[53:54]
	v_fma_f64 v[0:1], v[51:52], v[59:60], v[0:1]
	v_min_f64 v[45:46], v[45:46], 0
	v_mov_b32_e32 v49, 0
	v_dual_mov_b32 v50, 0x3ff00000 :: v_dual_mov_b32 v47, 0
	v_mov_b32_e32 v48, 0x408f4000
.LBB0_41:                               ;   Parent Loop BB0_11 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_cmp_eq_u32 s36, 1
	s_mov_b32 s16, s4
	s_cselect_b32 vcc_lo, -1, 0
	s_cmp_eq_u32 s36, 2
	v_cndmask_b32_e32 v51, v42, v3, vcc_lo
	v_cndmask_b32_e32 v53, v41, v2, vcc_lo
	s_cselect_b32 vcc_lo, -1, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_max_f64 v[47:48], v[47:48], v[47:48]
	s_add_u32 s36, s36, 1
	v_dual_cndmask_b32 v52, v51, v1 :: v_dual_cndmask_b32 v51, v53, v0
	s_addc_u32 s37, s37, 0
	s_cmp_lg_u32 s36, 3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_scale_f64 v[53:54], null, v[51:52], v[51:52], v[45:46]
	v_div_scale_f64 v[55:56], null, v[51:52], v[51:52], v[43:44]
	v_rcp_f64_e32 v[57:58], v[53:54]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_rcp_f64_e32 v[59:60], v[55:56]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[61:62], -v[53:54], v[57:58], 1.0
	v_fma_f64 v[63:64], -v[55:56], v[59:60], 1.0
	v_fma_f64 v[57:58], v[57:58], v[61:62], v[57:58]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_fma_f64 v[59:60], v[59:60], v[63:64], v[59:60]
	v_div_scale_f64 v[63:64], vcc_lo, v[45:46], v[51:52], v[45:46]
	v_fma_f64 v[61:62], -v[53:54], v[57:58], 1.0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[65:66], -v[55:56], v[59:60], 1.0
	v_fma_f64 v[57:58], v[57:58], v[61:62], v[57:58]
	v_div_scale_f64 v[61:62], s3, v[43:44], v[51:52], v[43:44]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[59:60], v[59:60], v[65:66], v[59:60]
	v_mul_f64 v[65:66], v[63:64], v[57:58]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f64 v[67:68], v[61:62], v[59:60]
	v_fma_f64 v[53:54], -v[53:54], v[65:66], v[63:64]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[55:56], -v[55:56], v[67:68], v[61:62]
	v_div_fmas_f64 v[53:54], v[53:54], v[57:58], v[65:66]
	s_mov_b32 vcc_lo, s3
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_div_fmas_f64 v[55:56], v[55:56], v[59:60], v[67:68]
	v_cmp_gt_f64_e32 vcc_lo, s[4:5], v[51:52]
	v_div_fixup_f64 v[53:54], v[53:54], v[51:52], v[45:46]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e32 v50, v50, v54, vcc_lo
	v_div_fixup_f64 v[55:56], v[55:56], v[51:52], v[43:44]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v49, v49, v53, vcc_lo
	v_cmp_lt_f64_e32 vcc_lo, s[16:17], v[51:52]
	v_dual_cndmask_b32 v50, v50, v56 :: v_dual_cndmask_b32 v49, v49, v55
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_max_f64 v[51:52], v[49:50], v[49:50]
	v_min_f64 v[47:48], v[51:52], v[47:48]
	s_cbranch_scc1 .LBB0_41
; %bb.42:                               ; %_Z14limit_gradientPdddd.exit726
                                        ;   in Loop: Header=BB0_11 Depth=1
	v_fma_f64 v[43:44], v[25:26], v[39:40], s[48:49]
	s_mov_b64 s[36:37], 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[43:44], v[43:44], v[47:48]
	v_min_f64 v[43:44], v[43:44], 1.0
	s_delay_alu instid0(VALU_DEP_1)
	v_mul_f64 v[40:41], v[41:42], v[43:44]
	v_mul_f64 v[2:3], v[2:3], v[43:44]
	v_mul_f64 v[0:1], v[0:1], v[43:44]
	v_dual_mov_b32 v43, v34 :: v_dual_mov_b32 v42, v33
.LBB0_43:                               ;   Parent Loop BB0_11 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b64 v[44:45], v[10:11], off
	s_cmp_eq_u32 s36, 1
	s_cselect_b32 vcc_lo, -1, 0
	s_cmp_eq_u32 s36, 2
	v_cndmask_b32_e32 v46, v41, v3, vcc_lo
	v_cndmask_b32_e32 v48, v40, v2, vcc_lo
	s_cselect_b32 vcc_lo, -1, 0
	s_add_u32 s36, s36, 1
	s_addc_u32 s37, s37, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v47, v46, v1 :: v_dual_cndmask_b32 v46, v48, v0
	s_cmp_eq_u32 s36, 3
	s_waitcnt vmcnt(0)
	v_add_f64 v[44:45], v[46:47], v[44:45]
	global_store_b64 v[42:43], v[44:45], off
	v_add_co_u32 v42, vcc_lo, v42, 8
	v_add_co_ci_u32_e32 v43, vcc_lo, 0, v43, vcc_lo
	s_cbranch_scc0 .LBB0_43
; %bb.44:                               ; %Flow1014
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_or_b32 s3, s57, exec_lo
                                        ; implicit-def: $vgpr58_vgpr59
.LBB0_45:                               ; %Flow1024
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_or_b32 exec_lo, exec_lo, s60
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 s16, s57, exec_lo
	s_and_b32 s3, s3, exec_lo
	s_and_b32 s59, s59, exec_lo
	s_or_b32 s57, s16, s3
.LBB0_46:                               ; %Flow1022
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_or_b32 exec_lo, exec_lo, s58
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 s3, s57, exec_lo
	s_and_b32 s16, s59, exec_lo
.LBB0_47:                               ; %Flow1019
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_and_not1_saveexec_b32 s36, s47
	s_cbranch_execz .LBB0_49
; %bb.48:                               ;   in Loop: Header=BB0_11 Depth=1
	global_load_b64 v[0:1], v[17:18], off
	s_or_b32 s16, s16, exec_lo
	s_waitcnt vmcnt(0)
	v_dual_mov_b32 v2, v0 :: v_dual_mov_b32 v3, v1
	global_store_b128 v[19:20], v[0:3], off
	global_load_b64 v[58:59], v[17:18], off
.LBB0_49:                               ; %Flow1025
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_or_b32 exec_lo, exec_lo, s36
	s_and_saveexec_b32 s36, s16
	s_cbranch_execz .LBB0_51
; %bb.50:                               ; %.loopexit766.sink.split
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_waitcnt vmcnt(0)
	global_store_b64 v[19:20], v[58:59], off offset:16
	global_load_b64 v[0:1], v[12:13], off
	s_or_b32 s3, s3, exec_lo
	s_waitcnt vmcnt(0)
	v_dual_mov_b32 v2, v0 :: v_dual_mov_b32 v3, v1
	global_store_b128 v[31:32], v[0:3], off
	global_load_b64 v[0:1], v[12:13], off
	s_waitcnt vmcnt(0)
	global_store_b64 v[31:32], v[0:1], off offset:16
	global_load_b64 v[0:1], v[10:11], off
	s_waitcnt vmcnt(0)
	v_dual_mov_b32 v2, v0 :: v_dual_mov_b32 v3, v1
	global_store_b128 v[33:34], v[0:3], off
	global_load_b64 v[0:1], v[10:11], off
	s_waitcnt vmcnt(0)
	global_store_b64 v[33:34], v[0:1], off offset:16
.LBB0_51:                               ; %Flow1026
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_or_b32 exec_lo, exec_lo, s36
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 s36, s46, exec_lo
	s_and_b32 s37, s56, exec_lo
	s_mov_b32 s16, -1
	s_or_b32 s46, s36, s37
	s_and_saveexec_b32 s36, s3
; %bb.52:                               ; %.loopexit766
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_and_not1_b32 s46, s46, exec_lo
	s_or_not1_b32 s16, s0, exec_lo
; %bb.53:                               ; %Flow1027
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_or_b32 exec_lo, exec_lo, s36
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s3, exec_lo, s16
	s_or_b32 s42, s3, s42
	s_and_not1_b32 s3, s43, exec_lo
	s_and_b32 s16, s46, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s43, s3, s16
	s_and_not1_b32 exec_lo, exec_lo, s42
	s_cbranch_execnz .LBB0_11
; %bb.54:                               ; %loop.exit.guard
	s_or_b32 exec_lo, exec_lo, s42
	s_xor_b32 s0, s43, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_saveexec_b32 s1, s0
	s_xor_b32 s1, exec_lo, s1
	s_cbranch_execz .LBB0_60
; %bb.55:
	v_dual_mov_b32 v17, 0 :: v_dual_add_nc_u32 v0, 2, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_mov_b32_e32 v1, v17
	v_lshlrev_b64 v[18:19], 3, v[16:17]
	v_mov_b32_e32 v9, v17
	v_lshlrev_b64 v[20:21], 3, v[0:1]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v6, vcc_lo, s18, v18
	v_add_co_ci_u32_e32 v7, vcc_lo, s19, v19, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v14, vcc_lo, s18, v20
	v_add_co_ci_u32_e32 v15, vcc_lo, s19, v21, vcc_lo
	v_lshlrev_b64 v[0:1], 3, v[8:9]
	s_clause 0x1
	global_load_b64 v[22:23], v[6:7], off
	global_load_b64 v[24:25], v[14:15], off
	v_add_co_u32 v2, vcc_lo, s18, v0
	v_add_co_ci_u32_e32 v3, vcc_lo, s19, v1, vcc_lo
	global_load_b64 v[8:9], v[2:3], off
	s_waitcnt vmcnt(1)
	v_add_f64 v[16:17], v[22:23], v[24:25]
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f64 v[8:9], v[16:17], -v[8:9]
	v_add_co_u32 v16, vcc_lo, s26, v0
	v_add_co_ci_u32_e32 v17, vcc_lo, s27, v1, vcc_lo
	global_store_b64 v[16:17], v[8:9], off
	s_clause 0x2
	global_load_b64 v[8:9], v[2:3], off
	global_load_b64 v[16:17], v[14:15], off
	global_load_b64 v[22:23], v[6:7], off
	s_waitcnt vmcnt(1)
	v_add_f64 v[8:9], v[8:9], v[16:17]
	v_add_co_u32 v16, vcc_lo, s26, v18
	v_add_co_ci_u32_e32 v17, vcc_lo, s27, v19, vcc_lo
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_3)
	v_add_f64 v[8:9], v[8:9], -v[22:23]
	global_store_b64 v[16:17], v[8:9], off
	s_clause 0x2
	global_load_b64 v[6:7], v[6:7], off
	global_load_b64 v[8:9], v[2:3], off
	global_load_b64 v[14:15], v[14:15], off
	s_waitcnt vmcnt(1)
	v_add_f64 v[6:7], v[8:9], v[6:7]
	v_add_co_u32 v8, vcc_lo, s22, v18
	v_add_co_ci_u32_e32 v9, vcc_lo, s23, v19, vcc_lo
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_3)
	v_add_f64 v[6:7], v[6:7], -v[14:15]
	v_add_co_u32 v14, vcc_lo, s26, v20
	v_add_co_ci_u32_e32 v15, vcc_lo, s27, v21, vcc_lo
	v_add_co_u32 v16, vcc_lo, s22, v20
	v_add_co_ci_u32_e32 v17, vcc_lo, s23, v21, vcc_lo
	global_store_b64 v[14:15], v[6:7], off
	s_clause 0x1
	global_load_b64 v[14:15], v[8:9], off
	global_load_b64 v[22:23], v[16:17], off
	v_add_co_u32 v6, vcc_lo, s22, v0
	v_add_co_ci_u32_e32 v7, vcc_lo, s23, v1, vcc_lo
	global_load_b64 v[24:25], v[6:7], off
	s_waitcnt vmcnt(1)
	v_add_f64 v[14:15], v[14:15], v[22:23]
	v_add_co_u32 v22, vcc_lo, s30, v0
	v_add_co_ci_u32_e32 v23, vcc_lo, s31, v1, vcc_lo
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_3)
	v_add_f64 v[14:15], v[14:15], -v[24:25]
	global_store_b64 v[22:23], v[14:15], off
	s_clause 0x2
	global_load_b64 v[14:15], v[6:7], off
	global_load_b64 v[22:23], v[16:17], off
	global_load_b64 v[24:25], v[8:9], off
	s_waitcnt vmcnt(1)
	v_add_f64 v[14:15], v[14:15], v[22:23]
	v_add_co_u32 v22, vcc_lo, s30, v18
	v_add_co_ci_u32_e32 v23, vcc_lo, s31, v19, vcc_lo
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_3)
	v_add_f64 v[14:15], v[14:15], -v[24:25]
	global_store_b64 v[22:23], v[14:15], off
	s_clause 0x2
	global_load_b64 v[8:9], v[8:9], off
	global_load_b64 v[14:15], v[6:7], off
	global_load_b64 v[16:17], v[16:17], off
	s_waitcnt vmcnt(1)
	v_add_f64 v[8:9], v[14:15], v[8:9]
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f64 v[8:9], v[8:9], -v[16:17]
	v_add_co_u32 v16, vcc_lo, s24, v18
	v_add_co_ci_u32_e32 v17, vcc_lo, s25, v19, vcc_lo
	v_add_co_u32 v14, vcc_lo, s30, v20
	v_add_co_ci_u32_e32 v15, vcc_lo, s31, v21, vcc_lo
	v_add_co_u32 v22, vcc_lo, s24, v20
	v_add_co_ci_u32_e32 v23, vcc_lo, s25, v21, vcc_lo
	global_store_b64 v[14:15], v[8:9], off
	s_clause 0x1
	global_load_b64 v[8:9], v[16:17], off
	global_load_b64 v[24:25], v[22:23], off
	v_add_co_u32 v14, vcc_lo, s24, v0
	v_add_co_ci_u32_e32 v15, vcc_lo, s25, v1, vcc_lo
	global_load_b64 v[26:27], v[14:15], off
	s_waitcnt vmcnt(1)
	v_add_f64 v[8:9], v[8:9], v[24:25]
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f64 v[24:25], v[8:9], -v[26:27]
	v_add_co_u32 v8, vcc_lo, s34, v0
	v_add_co_ci_u32_e32 v9, vcc_lo, s35, v1, vcc_lo
	v_add_co_u32 v18, vcc_lo, s34, v18
	v_add_co_ci_u32_e32 v19, vcc_lo, s35, v19, vcc_lo
	global_store_b64 v[8:9], v[24:25], off
	s_clause 0x2
	global_load_b64 v[24:25], v[14:15], off
	global_load_b64 v[26:27], v[22:23], off
	global_load_b64 v[28:29], v[16:17], off
	s_waitcnt vmcnt(1)
	v_add_f64 v[24:25], v[24:25], v[26:27]
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f64 v[24:25], v[24:25], -v[28:29]
	global_store_b64 v[18:19], v[24:25], off
	s_clause 0x2
	global_load_b64 v[16:17], v[16:17], off
	global_load_b64 v[18:19], v[14:15], off
	global_load_b64 v[22:23], v[22:23], off
	s_waitcnt vmcnt(1)
	v_add_f64 v[16:17], v[18:19], v[16:17]
	v_add_co_u32 v18, vcc_lo, s34, v20
	v_add_co_ci_u32_e32 v19, vcc_lo, s35, v21, vcc_lo
	s_and_not1_b32 vcc_lo, exec_lo, s33
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_3)
	v_add_f64 v[16:17], v[16:17], -v[22:23]
	global_store_b64 v[18:19], v[16:17], off
	s_cbranch_vccnz .LBB0_60
; %bb.56:
	v_add_co_u32 v16, vcc_lo, s52, v4
	v_add_co_ci_u32_e32 v17, vcc_lo, s53, v5, vcc_lo
	v_add_co_u32 v4, vcc_lo, s54, v4
	v_add_co_ci_u32_e32 v5, vcc_lo, s55, v5, vcc_lo
	global_load_b64 v[16:17], v[16:17], off
	s_mov_b64 s[0:1], 0
	s_waitcnt vmcnt(0)
	global_store_b64 v[12:13], v[16:17], off
	global_load_b64 v[12:13], v[4:5], off
	v_add_co_u32 v4, vcc_lo, s20, v0
	v_add_co_ci_u32_e32 v5, vcc_lo, s21, v1, vcc_lo
	s_waitcnt vmcnt(0)
	global_store_b64 v[10:11], v[12:13], off
.LBB0_57:                               ; =>This Inner Loop Header: Depth=1
	v_add_co_u32 v10, vcc_lo, v2, s0
	v_add_co_ci_u32_e32 v11, vcc_lo, s1, v3, vcc_lo
	v_add_co_u32 v12, vcc_lo, v4, s0
	v_add_co_ci_u32_e32 v13, vcc_lo, s1, v5, vcc_lo
	v_add_co_u32 v16, vcc_lo, v6, s0
	global_load_b64 v[10:11], v[10:11], off
	global_load_b64 v[12:13], v[12:13], off
	v_add_co_ci_u32_e32 v17, vcc_lo, s1, v7, vcc_lo
	global_load_b64 v[18:19], v[16:17], off
	s_waitcnt vmcnt(1)
	v_add_f64 v[10:11], v[10:11], -v[12:13]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_max_f64 v[10:11], v[10:11], 0
	s_waitcnt vmcnt(0)
	v_mul_f64 v[12:13], v[18:19], v[10:11]
	v_add_co_u32 v18, vcc_lo, v14, s0
	v_add_co_ci_u32_e32 v19, vcc_lo, s1, v15, vcc_lo
	s_add_u32 s0, s0, 8
	s_addc_u32 s1, s1, 0
	s_cmp_lg_u32 s0, 24
	global_store_b64 v[16:17], v[12:13], off
	global_load_b64 v[12:13], v[18:19], off
	s_waitcnt vmcnt(0)
	v_mul_f64 v[10:11], v[10:11], v[12:13]
	global_store_b64 v[18:19], v[10:11], off
	s_cbranch_scc1 .LBB0_57
; %bb.58:                               ; %.preheader.preheader
	v_add_co_u32 v2, vcc_lo, s26, v0
	v_add_co_ci_u32_e32 v3, vcc_lo, s27, v1, vcc_lo
	v_add_co_u32 v4, vcc_lo, s28, v0
	v_add_co_ci_u32_e32 v5, vcc_lo, s29, v1, vcc_lo
	v_add_co_u32 v0, vcc_lo, s30, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s31, v1, vcc_lo
	s_mov_b64 s[0:1], 0
.LBB0_59:                               ; %.preheader
                                        ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	v_add_co_u32 v6, vcc_lo, v2, s0
	v_add_co_ci_u32_e32 v7, vcc_lo, s1, v3, vcc_lo
	v_add_co_u32 v10, vcc_lo, v4, s0
	v_add_co_ci_u32_e32 v11, vcc_lo, s1, v5, vcc_lo
	v_add_co_u32 v12, vcc_lo, v0, s0
	global_load_b64 v[6:7], v[6:7], off
	global_load_b64 v[10:11], v[10:11], off
	v_add_co_ci_u32_e32 v13, vcc_lo, s1, v1, vcc_lo
	global_load_b64 v[14:15], v[12:13], off
	s_waitcnt vmcnt(1)
	v_add_f64 v[6:7], v[6:7], -v[10:11]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_max_f64 v[6:7], v[6:7], 0
	s_waitcnt vmcnt(0)
	v_mul_f64 v[10:11], v[14:15], v[6:7]
	v_add_co_u32 v14, vcc_lo, v8, s0
	v_add_co_ci_u32_e32 v15, vcc_lo, s1, v9, vcc_lo
	s_add_u32 s0, s0, 8
	s_addc_u32 s1, s1, 0
	s_cmp_lg_u32 s0, 24
	global_store_b64 v[12:13], v[10:11], off
	global_load_b64 v[10:11], v[14:15], off
	s_waitcnt vmcnt(0)
	v_mul_f64 v[6:7], v[6:7], v[10:11]
	global_store_b64 v[14:15], v[6:7], off
	s_cbranch_scc1 .LBB0_59
.LBB0_60:                               ; %.loopexit
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 512
		.amdhsa_user_sgpr_count 14
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 114
		.amdhsa_next_free_sgpr 61
		.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	_Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi, .Lfunc_end0-_Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 8736
; NumSgprs: 63
; NumVgprs: 114
; ScratchSize: 0
; MemoryBound: 1
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 7
; VGPRBlocks: 14
; NumSGPRsForWavesPerEU: 63
; NumVGPRsForWavesPerEU: 114
; Occupancy: 12
; WaveLimiterHint : 1
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .offset:         4
        .size:           4
        .value_kind:     by_value
      - .offset:         8
        .size:           4
        .value_kind:     by_value
      - .offset:         16
        .size:           8
        .value_kind:     by_value
      - .offset:         24
        .size:           8
        .value_kind:     by_value
      - .offset:         32
        .size:           8
        .value_kind:     by_value
      - .offset:         40
        .size:           8
        .value_kind:     by_value
      - .offset:         48
        .size:           8
        .value_kind:     by_value
      - .offset:         56
        .size:           8
        .value_kind:     by_value
      - .offset:         64
        .size:           8
        .value_kind:     by_value
      - .offset:         72
        .size:           8
        .value_kind:     by_value
      - .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
      - .address_space:  global
        .offset:         104
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         112
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         120
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         128
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         136
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         144
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         152
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         160
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         168
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         176
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         184
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         192
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         200
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         208
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         216
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         224
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         232
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         240
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         248
        .size:           8
        .value_kind:     global_buffer
      - .offset:         256
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         260
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         264
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         268
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         270
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         272
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         274
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         276
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         278
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         296
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         304
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         312
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         320
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 512
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi
    .private_segment_fixed_size: 0
    .sgpr_count:     63
    .sgpr_spill_count: 0
    .symbol:         _Z33_extrapolate_second_order_edge_swiiiddddddddPlS_PdS0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_S0_Pi.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     114
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
