gistfile1.txt

inline float __shfl_down_3(float v0, int v1, int v2) {
    local float mem[1024];
    int tid = get_local_id(0);
    int warpid = tid % 32;
    int warpstart = tid - warpid;
    mem[tid] = v0;
    //barrier(CLK_LOCAL_MEM_FENCE);
    int warpsrc = warpid + v1;
    warpsrc = warpsrc >= 32 ? warpid : warpsrc;
    return mem[warpstart + warpsrc];
}

inline float __shfl_down_2(float v0, int v1) {
    return __shfl_down_3(v0, v1, 32);
}

// based on https://community.amd.com/thread/167462
inline int __atomic_inc(global volatile int *ptr, int val) {
    while( true ){
        int old = *ptr;
        if( old >= val ) {
            if(old == atomic_cmpxchg(ptr, old, 0)) {
               break;
            }
        }
        else {
            if(old == atomic_cmpxchg(ptr, old, old+1)) {
                break;
            }
        }
    }
    return 0;
}

//float __atomic_add(global volatile float *ptr, float val);

inline float __atomic_add(global volatile float *ptr, float val) { // we need to actually implement this
    return 555;
}

struct class_Eigen__array {
    char f0[1];
};

struct Eigen__Sizes {
    char f0;
};

struct class_Eigen__array_0 {
    long f0;
};

struct class_Eigen__array_1 {
    long f0[1];
};

struct Eigen__DSizes {
    struct class_Eigen__array_1 f0;
};

struct class_Eigen__StreamInterface {
};

struct Eigen__GpuDevice {
    global struct class_Eigen__StreamInterface* f0;
    int f1;
    char f2[4];
};

struct class_Eigen__TensorMap {
    global float* f0;
    struct Eigen__DSizes f1;
};

struct Eigen__TensorEvaluator_2 {
    global float* f0;
    struct Eigen__DSizes f1;
    global struct Eigen__GpuDevice* f2;
    global struct class_Eigen__TensorMap* f3;
};

struct Eigen__internal__SumReducer {
    char f0;
};

struct Eigen__TensorEvaluator {
    struct class_Eigen__array f0;
    struct Eigen__Sizes f1;
    struct class_Eigen__array_0 f2;
    struct class_Eigen__array_1 f3;
    struct class_Eigen__array_1 f4;
    struct class_Eigen__array_1 f5;
    struct Eigen__TensorEvaluator_2 f6;
    struct Eigen__internal__SumReducer f7;
    global float* f8;
    global struct Eigen__GpuDevice* f9;
};

struct class_Eigen__array_nopointers {
    char f0[1];
};

struct Eigen__Sizes_nopointers {
    char f0;
};

struct class_Eigen__array_0_nopointers {
    long f0;
};

struct class_Eigen__array_1_nopointers {
    long f0[1];
};

struct Eigen__DSizes_nopointers {
    struct class_Eigen__array_1_nopointers f0;
};

struct Eigen__TensorEvaluator_2_nopointers {
        struct Eigen__DSizes_nopointers f1;
        };

struct Eigen__internal__SumReducer_nopointers {
    char f0;
};

struct Eigen__TensorEvaluator_nopointers {
    struct class_Eigen__array_nopointers f0;
    struct Eigen__Sizes_nopointers f1;
    struct class_Eigen__array_0_nopointers f2;
    struct class_Eigen__array_1_nopointers f3;
    struct class_Eigen__array_1_nopointers f4;
    struct class_Eigen__array_1_nopointers f5;
    struct Eigen__TensorEvaluator_2_nopointers f6;
    struct Eigen__internal__SumReducer_nopointers f7;
        };

struct Eigen__TensorEvaluator_5_nopointers {
        struct Eigen__DSizes_nopointers f1;
        };

struct Eigen__TensorEvaluator_4_nopointers {
    struct Eigen__TensorEvaluator_5_nopointers f0;
    struct Eigen__TensorEvaluator_nopointers f1;
};

struct Eigen__TensorEvaluator_5 {
    global float* f0;
    struct Eigen__DSizes f1;
    global struct Eigen__GpuDevice* f2;
    global struct class_Eigen__TensorMap* f3;
};

struct Eigen__TensorEvaluator_4 {
    struct Eigen__TensorEvaluator_5 f0;
    struct Eigen__TensorEvaluator f1;
};



float _ZN5Eigen8internal19InnerMostDimReducerINS_15TensorEvaluatorIKNS_17TensorReductionOpINS0_10SumReducerIfEEKNS_5arrayIlLm1EEEKNS_9TensorMapINS_6TensorIfLi1ELi0ElEELi0ENS_11MakePointerEEEEENS_9GpuDeviceEEES5_Lb1EE6reduceERKSI_llRS5_(struct Eigen__TensorEvaluator* self, long firstIndex, long numValuesToReduce, struct Eigen__internal__SumReducer* reducer);

kernel void _ZN5Eigen8internal19FullReductionKernelILi256ELi128ENS_15TensorEvaluatorIKNS_17TensorReductionOpINS0_10SumReducerIfEEKNS_5arrayIlLm1EEEKNS_9TensorMapINS_6TensorIfLi1ELi0ElEELi0ENS_11MakePointerEEEEENS_9GpuDeviceEEES5_lEEvT2_T1_T3_PNSK_15CoeffReturnTypeEPj(global struct Eigen__internal__SumReducer* reducer, long reducer_offset, global struct Eigen__TensorEvaluator_nopointers* input_nopointers, global float* input_ptr0, long input_ptr_offset0, global float* input_ptr1, long input_ptr_offset1, long num_coeffs, global float* output, long output_offset, global int* semaphore, long semaphore_offset) {
    float accum_05;
    float accum_05_prol;
    float accum_05_unr;
    float accum_13_ph;
    long i_06;
    long i_06_prol;
    long i_06_unr;
    long prol_iter;
    float v_lcssa;
    long v_lcssa10;
    float v_lcssa8;
    float v_lcssa9;
    float v_lcssa_unr;
    int v1;
    int v3;
    int v4;
    long v5;
    int v6;
    float v24;
    int v26;
    long v33;
    long v__i;
    global float* v56;
    float v42;
    float v44;
    float v46;
    float v48;
    long v61;
    long v63;
    long xtraiter;
    float v70;
    long v69;
    long prol_iter_sub;
    float v38;
    long v80;
    float v102;
    int v110;
    semaphore = (global int*)((global char *)semaphore + semaphore_offset);
    if(semaphore_offset == -1) {
        semaphore = 0;
    }
    output = (global float*)((global char *)output + output_offset);
    if(output_offset == -1) {
        output = 0;
    }
    input_ptr1 = (global float*)((global char *)input_ptr1 + input_ptr_offset1);
    if(input_ptr_offset1 == -1) {
        input_ptr1 = 0;
    }
    input_ptr0 = (global float*)((global char *)input_ptr0 + input_ptr_offset0);
    if(input_ptr_offset0 == -1) {
        input_ptr0 = 0;
    }
    reducer = (global struct Eigen__internal__SumReducer*)((global char *)reducer + reducer_offset);
    if(reducer_offset == -1) {
        reducer = 0;
    }
struct Eigen__TensorEvaluator input[1];
input[0].f0.f0[0] = input_nopointers[0].f0.f0[0];
input[0].f1.f0 = input_nopointers[0].f1.f0;
input[0].f2.f0 = input_nopointers[0].f2.f0;
input[0].f3.f0[0] = input_nopointers[0].f3.f0[0];
input[0].f4.f0[0] = input_nopointers[0].f4.f0[0];
input[0].f5.f0[0] = input_nopointers[0].f5.f0[0];
input[0].f6.f1.f0.f0[0] = input_nopointers[0].f6.f1.f0.f0[0];
input[0].f7.f0 = input_nopointers[0].f7.f0;
input[0].f6.f0 = input_ptr0;
input[0].f8 = input_ptr1;

    label0:;
    v1 = get_group_id(0);
    v3 = get_local_id(0);
    v4 = (v1 << 15) + v3;
    v5 = v4;
    v6 = get_num_groups(0);
    if(!(v6 == 1)) {
        goto v9;
    }
    v10:;
    if(!(v4 == 0)) {
        goto v_loopexit;
    }
    v13:;
    output[0] = 0.0f;
        goto v_loopexit;
    v9:;
    if(!(v3 == 0)) {
        goto v_loopexit;
    }
    v18:;
    if (atomic_cmpxchg(semaphore, 0, 1) == 0) {
        goto v22;
    }
    v_preheader2_preheader:;
        goto v_preheader2;
    v22:;
    v24 = atomic_xchg(output, 0.0f);
    barrier(CLK_GLOBAL_MEM_FENCE);
    v26 = atomic_xchg(semaphore, 2);
        goto v_loopexit;
    v_preheader2:;
    if (atomic_cmpxchg(semaphore, 2, 2) < 2) {
        goto v_preheader2;
    }
    v_loopexit_loopexit:;
    v_loopexit:;
    barrier(CLK_GLOBAL_MEM_FENCE);
    v33 = num_coeffs - v5;
    v__i = (v33 > 32768) ? 32768 : v33;
    if (v__i > 0) {
        goto v_lr_ph;
    } else {
        accum_13_ph = 0.0f;
        goto v_preheader_preheader;
    }
    label37:;
        v_lcssa = v_lcssa8;
    v_preheader_preheader_loopexit:;
        accum_13_ph = v_lcssa;
    v_preheader_preheader:;
    v42 = accum_13_ph + __shfl_down_3(accum_13_ph, 16, 32);
    v44 = v42 + __shfl_down_3(v42, 8, 32);
    v46 = v44 + __shfl_down_3(v44, 4, 32);
    v48 = v46 + __shfl_down_3(v46, 2, 32);
    if ((v3 & 31) == 0) {
        goto v54;
    } else {
        goto v53;
    }
    v_lr_ph:;
    v56 = (&input[0].f6.f0)[0];
    v61 = ((v3 + (v1 << 15)) + -1) - num_coeffs;
    v63 = -2 - ((v61 > -32769) ? v61 : -32769);
    xtraiter = ((v63 >> 8) + 1) & 3;
    if (xtraiter == 0) {
        i_06_unr = 0;
accum_05_unr = 0.0f;
        goto v_lr_ph_split;
    }
    v_preheader:;
        i_06_prol = 0;
accum_05_prol = 0.0f;
prol_iter = xtraiter;
    v68:;
    v70 = accum_05_prol + ((&v56[i_06_prol + v5])[0]);
    v69 = i_06_prol + 256;
    prol_iter_sub = prol_iter + -1;
    if (prol_iter_sub == 0) {
        v_lcssa10 = v69;
v_lcssa9 = v70;
    } else {
        i_06_prol = v69;
accum_05_prol = v70;
prol_iter = prol_iter_sub;
        goto v68;
    }
    v_lr_ph_split_loopexit:;
        i_06_unr = v_lcssa10;
accum_05_unr = v_lcssa9;
v_lcssa_unr = v_lcssa9;
    v_lr_ph_split:;
    if (v63 < 768) {
        v_lcssa = v_lcssa_unr;
        goto v_preheader_preheader_loopexit;
    }
    v_lr_ph_split_split:;
        i_06 = i_06_unr;
accum_05 = accum_05_unr;
    v79:;
    v38 = (((accum_05 + ((&v56[i_06 + v5])[0])) + ((&v56[(i_06 + 256) + v5])[0])) + ((&v56[(i_06 + 512) + v5])[0])) + ((&v56[(i_06 + 768) + v5])[0]);
    v80 = i_06 + 1024;
    if (v80 < v__i) {
        i_06 = v80;
accum_05 = v38;
        goto v79;
    } else {
        v_lcssa8 = v38;
        goto label37;
    }
    v54:;
    v102 = __atomic_add(output, (v48 + __shfl_down_3(v48, 1, 32)));
    v53:;
    if(!((v3 == 0) & (v6 > 1))) {
        goto v107;
    }
    v108:;
    v110 = __atomic_inc(semaphore, (v6 + 1));
    v107:;
    return;
}

kernel void _ZN5Eigen8internal19ReductionInitKernelIflEEvT_T0_PS2_(float val, long num_preserved_coeffs, global float* output, long output_offset) {
    long i_01;
    long v6;
    int v2;
    long v16;
    output = (global float*)((global char *)output + output_offset);
    if(output_offset == -1) {
        output = 0;
    }

    label0:;
    v2 = get_local_size(0);
    v6 = (v2 * get_group_id(0)) + get_local_id(0);
    if(!(v6 < num_preserved_coeffs)) {
        goto v__crit_edge;
    }
    v_lr_ph_preheader:;
        i_01 = v6;
    goto v_lr_ph;
    label13:;
    v__crit_edge:;
    return;
    v_lr_ph:;
    (&output[i_01])[0] = val;
    v16 = i_01 + (get_num_groups(0) * v2);
    if (v16 < num_preserved_coeffs) {
        i_01 = v16;
        goto v_lr_ph;
    } else {
        goto label13;
    }
}

kernel void _ZN5Eigen8internal20InnerReductionKernelILi128ENS_15TensorEvaluatorIKNS_17TensorReductionOpINS0_10SumReducerIfEEKNS_5arrayIlLm1EEEKNS_9TensorMapINS_6TensorIfLi1ELi0ElEELi0ENS_11MakePointerEEEEENS_9GpuDeviceEEES5_lEEvT1_T0_T2_SL_PNSK_15CoeffReturnTypeE(global struct Eigen__internal__SumReducer* reducer, long reducer_offset, global struct Eigen__TensorEvaluator_nopointers* input_nopointers, global float* input_ptr0, long input_ptr_offset0, global float* input_ptr1, long input_ptr_offset1, long num_coeffs_to_reduce, long num_preserved_coeffs, global float* output, long output_offset) {
    long col_011;
    long i1_014;
    long i_016;
    long j_09;
    float reduced_val_08;
    float reduced_val_08_lcssa;
    float reduced_val_110;
    float reduced_val_312_ph;
    long v_lcssa;
    float v_lcssa40;
    float v_lcssa41;
    long v19;
    float v179;
    float v181;
    float v183;
    float v185;
    long v3;
    int v1;
    long v6;
    long v7;
    int v8;
    int v11;
    int v12;
    long v27;
    long v32;
    long v35;
    global float* v39;
    long v49;
    long v56;
    long v57;
    long v58;
    long v61;
    long v67;
    float v62;
    long v165;
    float v170;
    long v169;
    float v191;
    long v48;
    output = (global float*)((global char *)output + output_offset);
    if(output_offset == -1) {
        output = 0;
    }
    input_ptr1 = (global float*)((global char *)input_ptr1 + input_ptr_offset1);
    if(input_ptr_offset1 == -1) {
        input_ptr1 = 0;
    }
    input_ptr0 = (global float*)((global char *)input_ptr0 + input_ptr_offset0);
    if(input_ptr_offset0 == -1) {
        input_ptr0 = 0;
    }
    reducer = (global struct Eigen__internal__SumReducer*)((global char *)reducer + reducer_offset);
    if(reducer_offset == -1) {
        reducer = 0;
    }
struct Eigen__TensorEvaluator input[1];
input[0].f0.f0[0] = input_nopointers[0].f0.f0[0];
input[0].f1.f0 = input_nopointers[0].f1.f0;
input[0].f2.f0 = input_nopointers[0].f2.f0;
input[0].f3.f0[0] = input_nopointers[0].f3.f0[0];
input[0].f4.f0[0] = input_nopointers[0].f4.f0[0];
input[0].f5.f0[0] = input_nopointers[0].f5.f0[0];
input[0].f6.f1.f0.f0[0] = input_nopointers[0].f6.f1.f0.f0[0];
input[0].f7.f0 = input_nopointers[0].f7.f0;
input[0].f6.f0 = input_ptr0;
input[0].f8 = input_ptr1;

    label0:;
    v1 = get_local_size(0);
    v3 = v1 << 7;
    v6 = ((num_coeffs_to_reduce + -1) + v3) / v3;
    v7 = v6 * num_preserved_coeffs;
    v8 = get_num_groups(0);
    v11 = get_group_id(0);
    v12 = get_local_id(0);
    if(!(v8 == 1)) {
        goto v15;
    }
    v16:;
    v19 = (v11 * v1) + v12;
    if(!(v19 < num_preserved_coeffs)) {
        goto v__crit_edge19;
    }
    v_lr_ph18_preheader:;
        i_016 = v19;
    goto v_lr_ph18;
    label23:;
    v__crit_edge19:;
    barrier(CLK_GLOBAL_MEM_FENCE);
        goto v15;
    v_lr_ph18:;
    (&output[i_016])[0] = 0.0f;
    v27 = i_016 + (v8 * v1);
    if (v27 < num_preserved_coeffs) {
        i_016 = v27;
        goto v_lr_ph18;
    } else {
        goto label23;
    }
    v15:;
    v32 = v11;
    if(!(v32 < v7)) {
        goto v__crit_edge;
    }
    v_lr_ph15:;
    v35 = v1;
    v39 = (&input[0].f6.f0)[0];
        i1_014 = v32;
    goto v44;
    label45:;
    v__crit_edge:;
    return;
    v44:;
    v49 = i1_014 / v6;
    if(!(v49 < num_preserved_coeffs)) {
        goto v52;
    }
    v53:;
    v56 = ((v35 << 7) * (i1_014 % v6)) + v12;
    v57 = v49 * num_coeffs_to_reduce;
    v58 = v56 + v57;
        j_09 = 0;
reduced_val_08 = 0.0f;
    v60:;
    v61 = j_09 + 16;
    v67 = j_09 * v35;
    if(!(((j_09 | 15) * v35) + v56 < num_coeffs_to_reduce)) {
        v_lcssa = v67;
reduced_val_08_lcssa = reduced_val_08;
        goto v69;
    }
    v_preheader_preheader:;
    v62 = (((((((((((((((reduced_val_08 + ((&v39[v58 + v67])[0])) + ((&v39[v58 + ((j_09 | 1) * v35)])[0])) + ((&v39[v58 + ((j_09 | 2) * v35)])[0])) + ((&v39[v58 + ((j_09 | 3) * v35)])[0])) + ((&v39[v58 + ((j_09 | 4) * v35)])[0])) + ((&v39[v58 + ((j_09 | 5) * v35)])[0])) + ((&v39[v58 + ((j_09 | 6) * v35)])[0])) + ((&v39[v58 + ((j_09 | 7) * v35)])[0])) + ((&v39[v58 + ((j_09 | 8) * v35)])[0])) + ((&v39[v58 + ((j_09 | 9) * v35)])[0])) + ((&v39[v58 + ((j_09 | 10) * v35)])[0])) + ((&v39[v58 + ((j_09 | 11) * v35)])[0])) + ((&v39[v58 + ((j_09 | 12) * v35)])[0])) + ((&v39[v58 + ((j_09 | 13) * v35)])[0])) + ((&v39[v58 + ((j_09 | 14) * v35)])[0])) + ((&v39[v58 + ((j_09 | 15) * v35)])[0]);
    if (v61 < 128) {
        j_09 = v61;
reduced_val_08 = v62;
        goto v60;
    } else {
        v_lcssa40 = v62;
        goto v_critedge_preheader_loopexit33;
    }
    v69:;
    v165 = v_lcssa + v56;
    if(!(v165 < num_coeffs_to_reduce)) {
        reduced_val_312_ph = reduced_val_08_lcssa;
        goto v_critedge_preheader;
    }
    v_lr_ph_preheader:;
        col_011 = v165;
reduced_val_110 = reduced_val_08_lcssa;
    v_lr_ph:;
    v170 = reduced_val_110 + ((&v39[col_011 + v57])[0]);
    v169 = col_011 + v35;
    if (v169 < num_coeffs_to_reduce) {
        col_011 = v169;
reduced_val_110 = v170;
        goto v_lr_ph;
    } else {
        v_lcssa41 = v170;
    }
    v_critedge_preheader_loopexit:;
        reduced_val_312_ph = v_lcssa41;
    goto v_critedge_preheader;
    v_critedge_preheader_loopexit33:;
        reduced_val_312_ph = v_lcssa40;
    v_critedge_preheader:;
    v179 = reduced_val_312_ph + __shfl_down_2(reduced_val_312_ph, 16);
    v181 = v179 + __shfl_down_2(v179, 8);
    v183 = v181 + __shfl_down_2(v181, 4);
    v185 = v183 + __shfl_down_2(v183, 2);
    if(!((v12 & 31) == 0)) {
        goto v52;
    }
    v188:;
    v191 = __atomic_add((&output[v49]), (v185 + __shfl_down_2(v185, 1)));
    v52:;
    v48 = i1_014 + v8;
    if (v48 < v7) {
        i1_014 = v48;
        goto v44;
    } else {
        goto label45;
    }
}

kernel void _ZN5Eigen8internal20OuterReductionKernelILi16ENS_15TensorEvaluatorIKNS_17TensorReductionOpINS0_10SumReducerIfEEKNS_5arrayIlLm1EEEKNS_9TensorMapINS_6TensorIfLi1ELi0ElEELi0ENS_11MakePointerEEEEENS_9GpuDeviceEEES5_lEEvT1_T0_T2_SL_PNSK_15CoeffReturnTypeE(global struct Eigen__internal__SumReducer* reducer, long reducer_offset, global struct Eigen__TensorEvaluator_nopointers* input_nopointers, global float* input_ptr0, long input_ptr_offset0, global float* input_ptr1, long input_ptr_offset1, long num_coeffs_to_reduce, long num_preserved_coeffs, global float* output, long output_offset) {
    long i1_04;
    long i_08;
    long j_03;
    float reduced_val_02;
    float reduced_val_0_lcssa;
    float v_lcssa;
    int v1;
    int v2;
    long v4;
    long v9;
    long v20;
    long v27;
    long v38;
    long v40;
    long v41;
    long v__i;
    float v50;
    long v37;
    float v47;
    long v53;
    output = (global float*)((global char *)output + output_offset);
    if(output_offset == -1) {
        output = 0;
    }
    input_ptr1 = (global float*)((global char *)input_ptr1 + input_ptr_offset1);
    if(input_ptr_offset1 == -1) {
        input_ptr1 = 0;
    }
    input_ptr0 = (global float*)((global char *)input_ptr0 + input_ptr_offset0);
    if(input_ptr_offset0 == -1) {
        input_ptr0 = 0;
    }
    reducer = (global struct Eigen__internal__SumReducer*)((global char *)reducer + reducer_offset);
    if(reducer_offset == -1) {
        reducer = 0;
    }
struct Eigen__TensorEvaluator input[1];
input[0].f0.f0[0] = input_nopointers[0].f0.f0[0];
input[0].f1.f0 = input_nopointers[0].f1.f0;
input[0].f2.f0 = input_nopointers[0].f2.f0;
input[0].f3.f0[0] = input_nopointers[0].f3.f0[0];
input[0].f4.f0[0] = input_nopointers[0].f4.f0[0];
input[0].f5.f0[0] = input_nopointers[0].f5.f0[0];
input[0].f6.f1.f0.f0[0] = input_nopointers[0].f6.f1.f0.f0[0];
input[0].f7.f0 = input_nopointers[0].f7.f0;
input[0].f6.f0 = input_ptr0;
input[0].f8 = input_ptr1;

    label0:;
    v1 = get_local_size(0);
    v2 = get_num_groups(0);
    v4 = v2 * v1;
    v9 = (get_group_id(0) * v1) + get_local_id(0);
    if(!(v2 == 1)) {
        goto v12;
    }
    v_preheader:;
    if(!(v9 < num_preserved_coeffs)) {
        goto v__crit_edge10;
    }
    v_lr_ph9_preheader:;
        i_08 = v9;
    goto v_lr_ph9;
    label16:;
    v__crit_edge10:;
    barrier(CLK_GLOBAL_MEM_FENCE);
        goto v12;
    v_lr_ph9:;
    (&output[i_08])[0] = 0.0f;
    v20 = i_08 + v4;
    if (v20 < num_preserved_coeffs) {
        i_08 = v20;
        goto v_lr_ph9;
    } else {
        goto label16;
    }
    v12:;
    v27 = ((num_coeffs_to_reduce + 15) / 16) * num_preserved_coeffs;
    if(!(v9 < v27)) {
        goto v__crit_edge7;
    }
    v_lr_ph6:;
        i1_04 = v9;
    goto v33;
    label34:;
    v__crit_edge7:;
    return;
    v33:;
    v38 = i1_04 % num_preserved_coeffs;
    v40 = (i1_04 / num_preserved_coeffs) << 4;
    v41 = v40 + 16;
    v__i = (v41 > num_coeffs_to_reduce) ? num_coeffs_to_reduce : v41;
    if(!(v40 < v__i)) {
        reduced_val_0_lcssa = 0.0f;
        goto v__crit_edge;
    }
    v_lr_ph_preheader:;
        j_03 = v40;
reduced_val_02 = 0.0f;
    goto v_lr_ph;
    label46:;
        reduced_val_0_lcssa = v_lcssa;
    v__crit_edge:;
    v50 = __atomic_add((&output[v38]), reduced_val_0_lcssa);
    v37 = i1_04 + v4;
    if (v37 < v27) {
        i1_04 = v37;
        goto v33;
    } else {
        goto label34;
    }
    v_lr_ph:;
    v47 = reduced_val_02 + ((&((&input[0].f6.f0)[0])[(j_03 * num_preserved_coeffs) + v38])[0]);
    v53 = j_03 + 1;
    if (v53 < v__i) {
        j_03 = v53;
reduced_val_02 = v47;
        goto v_lr_ph;
    } else {
        v_lcssa = v47;
        goto label46;
    }
}

kernel void _ZN5Eigen8internal15EigenMetaKernelINS_15TensorEvaluatorIKNS_14TensorAssignOpINS_9TensorMapINS_6TensorIfLi1ELi0ElEELi0ENS_11MakePointerEEEKNS_17TensorReductionOpINS0_10SumReducerIfEEKNS_5arrayIlLm1EEEKS8_EEEENS_9GpuDeviceEEElEEvT_T0_(global struct Eigen__TensorEvaluator_4_nopointers* eval_nopointers, global float* eval_ptr0, long eval_ptr_offset0, global float* eval_ptr1, long eval_ptr_offset1, global float* eval_ptr2, long eval_ptr_offset2, long size) {
    long i_01_i;
    float v_0_i_i_i;
    int v2;
    long v6;
    global float* v_pre;
    char* v14;
    float v25;
    long v28;
    float v30;
    long v20;
    eval_ptr2 = (global float*)((global char *)eval_ptr2 + eval_ptr_offset2);
    if(eval_ptr_offset2 == -1) {
        eval_ptr2 = 0;
    }
    eval_ptr1 = (global float*)((global char *)eval_ptr1 + eval_ptr_offset1);
    if(eval_ptr_offset1 == -1) {
        eval_ptr1 = 0;
    }
    eval_ptr0 = (global float*)((global char *)eval_ptr0 + eval_ptr_offset0);
    if(eval_ptr_offset0 == -1) {
        eval_ptr0 = 0;
    }
struct Eigen__TensorEvaluator_4 eval[1];
eval[0].f0.f1.f0.f0[0] = eval_nopointers[0].f0.f1.f0.f0[0];
eval[0].f1.f0.f0[0] = eval_nopointers[0].f1.f0.f0[0];
eval[0].f1.f1.f0 = eval_nopointers[0].f1.f1.f0;
eval[0].f1.f2.f0 = eval_nopointers[0].f1.f2.f0;
eval[0].f1.f3.f0[0] = eval_nopointers[0].f1.f3.f0[0];
eval[0].f1.f4.f0[0] = eval_nopointers[0].f1.f4.f0[0];
eval[0].f1.f5.f0[0] = eval_nopointers[0].f1.f5.f0[0];
eval[0].f1.f6.f1.f0.f0[0] = eval_nopointers[0].f1.f6.f1.f0.f0[0];
eval[0].f1.f7.f0 = eval_nopointers[0].f1.f7.f0;
eval[0].f0.f0 = eval_ptr0;
eval[0].f1.f6.f0 = eval_ptr1;
eval[0].f1.f8 = eval_ptr2;

    label0:;
    struct Eigen__internal__SumReducer reducer_i_i_i[1];
;
    v2 = get_local_size(0);
    v6 = (v2 * get_group_id(0)) + get_local_id(0);
    if(!(v6 < size)) {
        goto _ZN5Eigen8internal19EigenMetaKernelEvalINS_15TensorEvaluatorIKNS_14TensorAssignOpINS_9TensorMapINS_6TensorIfLi1ELi0ElEELi0ENS_11MakePointerEEEKNS_17TensorReductionOpINS0_10SumReducerIfEEKNS_5arrayIlLm1EEEKS8_EEEENS_9GpuDeviceEEElLb0EE3runERSL_lll_exit;
    }
    v_lr_ph_i:;
    v14 = &reducer_i_i_i[0].f0;
    v_pre = (&eval[0].f1.f8)[0];
        i_01_i = v6;
    v19:;
    if (v_pre == 0) {
        goto v23;
    }
    v22:;
    v25 = (&v_pre[i_01_i])[0];
        v_0_i_i_i = v25;
    goto _ZN5Eigen15TensorEvaluatorIKNS_14TensorAssignOpINS_9TensorMapINS_6TensorIfLi1ELi0ElEELi0ENS_11MakePointerEEEKNS_17TensorReductionOpINS_8internal10SumReducerIfEEKNS_5arrayIlLm1EEEKS6_EEEENS_9GpuDeviceEE10evalScalarEl_exit_i;
    v23:;
    ;
    v28 = (&eval[0].f1.f3.f0[0])[0];
    v30 = _ZN5Eigen8internal19InnerMostDimReducerINS_15TensorEvaluatorIKNS_17TensorReductionOpINS0_10SumReducerIfEEKNS_5arrayIlLm1EEEKNS_9TensorMapINS_6TensorIfLi1ELi0ElEELi0ENS_11MakePointerEEEEENS_9GpuDeviceEEES5_Lb1EE6reduceERKSI_llRS5_((&eval[0].f1), (v28 * i_01_i), v28, reducer_i_i_i);
    ;
        v_0_i_i_i = v30;
    _ZN5Eigen15TensorEvaluatorIKNS_14TensorAssignOpINS_9TensorMapINS_6TensorIfLi1ELi0ElEELi0ENS_11MakePointerEEEKNS_17TensorReductionOpINS_8internal10SumReducerIfEEKNS_5arrayIlLm1EEEKS6_EEEENS_9GpuDeviceEE10evalScalarEl_exit_i:;
    (&((&eval[0].f0.f0)[0])[i_01_i])[0] = v_0_i_i_i;
    v20 = i_01_i + (get_num_groups(0) * v2);
    if (v20 < size) {
        i_01_i = v20;
        goto v19;
    }
    _ZN5Eigen8internal19EigenMetaKernelEvalINS_15TensorEvaluatorIKNS_14TensorAssignOpINS_9TensorMapINS_6TensorIfLi1ELi0ElEELi0ENS_11MakePointerEEEKNS_17TensorReductionOpINS0_10SumReducerIfEEKNS_5arrayIlLm1EEEKS8_EEEENS_9GpuDeviceEEElLb0EE3runERSL_lll_exit_loopexit:;
    _ZN5Eigen8internal19EigenMetaKernelEvalINS_15TensorEvaluatorIKNS_14TensorAssignOpINS_9TensorMapINS_6TensorIfLi1ELi0ElEELi0ENS_11MakePointerEEEKNS_17TensorReductionOpINS0_10SumReducerIfEEKNS_5arrayIlLm1EEEKS8_EEEENS_9GpuDeviceEEElLb0EE3runERSL_lll_exit:;
    return;
}

float _ZN5Eigen8internal19InnerMostDimReducerINS_15TensorEvaluatorIKNS_17TensorReductionOpINS0_10SumReducerIfEEKNS_5arrayIlLm1EEEKNS_9TensorMapINS_6TensorIfLi1ELi0ElEELi0ENS_11MakePointerEEEEENS_9GpuDeviceEEES5_Lb1EE6reduceERKSI_llRS5_(struct Eigen__TensorEvaluator* self, long firstIndex, long numValuesToReduce, struct Eigen__internal__SumReducer* reducer) {
    float accum_02;
    float accum_02_prol;
    float accum_02_unr;
    float accum_0_lcssa;
    long j1_03;
    long j1_03_prol;
    long j1_03_unr;
    long j_08;
    float p_sroa_0_04;
    float p_sroa_0_0_lcssa;
    float p_sroa_10_06;
    float p_sroa_10_0_lcssa;
    float p_sroa_14_07;
    float p_sroa_14_0_lcssa;
    float p_sroa_6_05;
    float p_sroa_6_0_lcssa;
    long prol_iter;
    float v_lcssa;
    float v_lcssa32;
    float v_lcssa33;
    float v_lcssa34;
    float v_lcssa35;
    float v_lcssa36;
    float v_lcssa37;
    long v_lcssa38;
    float v_lcssa_unr;
    long v1;
    long v2;
    float4 v3;
    float v4;
    float v5;
    float v6;
    float v7;
    long xtraiter;
    global float* v22;
    long v23;
    float v31;
    long v30;
    long prol_iter_sub;
    global float* v44;
    float4 v52;
    float4 v61;
    float v17;
    float v16;
    float v15;
    float v14;
    long v41;
    float v65;
    long v73;
    label0:;
    v1 = numValuesToReduce / 4;
    v2 = v1 << 2;
    v3 = (float4)(0.0f, 0.0f, 0.0f, 0.0f);
    v4 = ((float*)&v3)[0];
    v5 = ((float*)&v3)[1];
    v6 = ((float*)&v3)[2];
    v7 = ((float*)&v3)[3];
    if(!(numValuesToReduce > 3)) {
        p_sroa_14_0_lcssa = v7;
p_sroa_10_0_lcssa = v6;
p_sroa_6_0_lcssa = v5;
p_sroa_0_0_lcssa = v4;
        goto v_preheader;
    }
    v_lr_ph10:;
        j_08 = 0;
p_sroa_14_07 = v7;
p_sroa_10_06 = v6;
p_sroa_6_05 = v5;
p_sroa_0_04 = v4;
    goto v12;
    label13:;
        p_sroa_14_0_lcssa = v_lcssa36;
p_sroa_10_0_lcssa = v_lcssa35;
p_sroa_6_0_lcssa = v_lcssa34;
p_sroa_0_0_lcssa = v_lcssa33;
    v_preheader:;
    if(!(v2 < numValuesToReduce)) {
        accum_0_lcssa = 0.0f;
        goto v__crit_edge;
    }
    v_lr_ph:;
    v22 = (&self[0].f6.f0)[0];
    v23 = v1 << 2;
    xtraiter = (numValuesToReduce - v23) & 3;
    if (xtraiter == 0) {
        j1_03_unr = v2;
accum_02_unr = 0.0f;
        goto v_lr_ph_split;
    }
    v_preheader31:;
        j1_03_prol = v2;
accum_02_prol = 0.0f;
prol_iter = xtraiter;
    v29:;
    v31 = accum_02_prol + ((&v22[j1_03_prol + firstIndex])[0]);
    v30 = j1_03_prol + 1;
    prol_iter_sub = prol_iter + -1;
    if (prol_iter_sub == 0) {
        v_lcssa38 = v30;
v_lcssa37 = v31;
    } else {
        j1_03_prol = v30;
accum_02_prol = v31;
prol_iter = prol_iter_sub;
        goto v29;
    }
    v_lr_ph_split_loopexit:;
        v_lcssa_unr = v_lcssa37;
j1_03_unr = v_lcssa38;
accum_02_unr = v_lcssa37;
    v_lr_ph_split:;
    if ((numValuesToReduce + -1) - v23 < 3) {
        v_lcssa = v_lcssa_unr;
        goto v__crit_edge_loopexit;
    }
    v_lr_ph_split_split:;
        j1_03 = j1_03_unr;
accum_02 = accum_02_unr;
    goto v40;
    v12:;
    v44 = &((&self[0].f6.f0)[0])[j_08 + firstIndex];
    v52 = (float4)(v44[0], ((&v44[1])[0]), ((&v44[2])[0]), ((&v44[3])[0]));
    v61 = (float4)((p_sroa_0_04 + (((float*)&v52)[0])), (p_sroa_6_05 + (((float*)&v52)[1])), (p_sroa_10_06 + (((float*)&v52)[2])), (p_sroa_14_07 + (((float*)&v52)[3])));
    v17 = ((float*)&v61)[0];
    v16 = ((float*)&v61)[1];
    v15 = ((float*)&v61)[2];
    v14 = ((float*)&v61)[3];
    v41 = j_08 + 4;
    if (v41 < v2) {
        j_08 = v41;
p_sroa_14_07 = v14;
p_sroa_10_06 = v15;
p_sroa_6_05 = v16;
p_sroa_0_04 = v17;
        goto v12;
    } else {
        v_lcssa36 = v14;
v_lcssa35 = v15;
v_lcssa34 = v16;
v_lcssa33 = v17;
        goto label13;
    }
    label64:;
        v_lcssa = v_lcssa32;
    v__crit_edge_loopexit:;
        accum_0_lcssa = v_lcssa;
    v__crit_edge:;
    return ((((p_sroa_0_0_lcssa + p_sroa_6_0_lcssa) + p_sroa_10_0_lcssa) + p_sroa_14_0_lcssa) + accum_0_lcssa);
    v40:;
    v65 = (((accum_02 + ((&v22[j1_03 + firstIndex])[0])) + ((&v22[(j1_03 + 1) + firstIndex])[0])) + ((&v22[(j1_03 + 2) + firstIndex])[0])) + ((&v22[(j1_03 + 3) + firstIndex])[0]);
    v73 = j1_03 + 4;
    if (v73 == numValuesToReduce) {
        v_lcssa32 = v65;
        goto label64;
    } else {
        j1_03 = v73;
accum_02 = v65;
        goto v40;
    }
}