extern "C" {
typedef struct {
    int e0;
    char* e1;
} struct_Buffer_6327;
typedef struct {
    struct_Buffer_6327 e0;
    struct_Buffer_6327 e1;
    int e2;
    int e3;
} struct_image_6326;
typedef struct {
    struct_Buffer_6327 e0;
    int e1;
    int e2;
} struct_filter_6332;

__device__ inline int threadIdx_x() { return threadIdx.x; }
__device__ inline int threadIdx_y() { return threadIdx.y; }
__device__ inline int threadIdx_z() { return threadIdx.z; }
__device__ inline int blockIdx_x() { return blockIdx.x; }
__device__ inline int blockIdx_y() { return blockIdx.y; }
__device__ inline int blockIdx_z() { return blockIdx.z; }
__device__ inline int blockDim_x() { return blockDim.x; }
__device__ inline int blockDim_y() { return blockDim.y; }
__device__ inline int blockDim_z() { return blockDim.z; }
__device__ inline int gridDim_x() { return gridDim.x; }
__device__ inline int gridDim_y() { return gridDim.y; }
__device__ inline int gridDim_z() { return gridDim.z; }
__global__ void lambda_25292(struct_image_6326, struct_Buffer_6327, double*, struct_filter_6332);
__global__ void lambda_25546(struct_image_6326, struct_filter_6332, double*, double*, struct_Buffer_6327);

__global__ __launch_bounds__ (128 * 1 * 1) void lambda_25292(struct_image_6326 _25295_30003, struct_Buffer_6327 _25296_30004, double* _25297_30005, struct_filter_6332 _25298_30006) {
    __shared__ double ds_img[134][7];
    int  bdimx_30012;
    int pbdimx_30012;
    int  bdimy_30018;
    int pbdimy_30018;
    int  bidx_30024;
    int pbidx_30024;
    int  bidy_30030;
    int pbidy_30030;
    int  tidx_30036;
    int ptidx_30036;
    int  tidy_30042;
    int ptidy_30042;
    double*  reserve_shared_30050;
    double* preserve_shared_30050;
    int  _30061;
    int p_30061;
    int  _30117;
    int p_30117;
    int  _30183;
    int p_30183;
    double  sum_30185;
    double psum_30185;
    int  _30122;
    int p_30122;
    #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
    bdimx_30012 = blockDim_x();
    pbdimx_30012 = bdimx_30012;
    l30010: ;
        bdimx_30012 = pbdimx_30012;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        bdimy_30018 = blockDim_y();
        pbdimy_30018 = bdimy_30018;
    l30016: ;
        bdimy_30018 = pbdimy_30018;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        bidx_30024 = blockIdx_x();
        pbidx_30024 = bidx_30024;
    l30022: ;
        bidx_30024 = pbidx_30024;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        bidy_30030 = blockIdx_y();
        pbidy_30030 = bidy_30030;
    l30028: ;
        bidy_30030 = pbidy_30030;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        tidx_30036 = threadIdx_x();
        ptidx_30036 = tidx_30036;
    l30034: ;
        tidx_30036 = ptidx_30036;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        tidy_30042 = threadIdx_y();
        ptidy_30042 = tidy_30042;
    l30040: ;
        tidy_30042 = ptidy_30042;
        #line 201 "gpu_device_shm.impala"
        __shared__ double reserver_reserve_shared_30050[938];
        preserve_shared_30050 = reserver_reserve_shared_30050;
    l30048: ;
        reserve_shared_30050 = preserve_shared_30050;
        #line 193 "gpu_device_shm.impala"
        int _30054;
        _30054 = _25298_30006.e2;
        #line 189 "gpu_device_shm.impala"
        int _30084;
        _30084 = bidy_30030 * bdimy_30018;
        #line 11 "main.impala"
        int _30081;
        _30081 = _25295_30003.e2;
        #line 187 "gpu_device_shm.impala"
        int _30075;
        _30075 = bidx_30024 * bdimx_30012;
        #line 11 "main.impala"
        int _30090;
        _30090 = _25295_30003.e3;
        #line 4 "gaussian.impala"
        int _30064;
        _30064 = _25298_30006.e1;
        #line 193 "gpu_device_shm.impala"
        int extend_height_30055;
        extend_height_30055 = _30054 / 2;
        #line 189 "gpu_device_shm.impala"
        int gid_y_30085;
        gid_y_30085 = _30084 + tidy_30042;
        #line 187 "gpu_device_shm.impala"
        int gid_x_30076;
        gid_x_30076 = _30075 + tidx_30036;
        #line 4 "gaussian.impala"
        int h_anchor_30065;
        h_anchor_30065 = _30064 / 2;
        #line 197 "gpu_device_shm.impala"
        int _30056;
        _30056 = 2 * extend_height_30055;
        #line 195 "gpu_device_shm.impala"
        int _30066;
        _30066 = 2 * h_anchor_30065;
        #line 197 "gpu_device_shm.impala"
        int shm_dimy_30057;
        shm_dimy_30057 = bdimy_30018 + _30056;
        #line 195 "gpu_device_shm.impala"
        int shm_dimx_30067;
        shm_dimx_30067 = bdimx_30012 + _30066;
        #line 38 "gpu_device_shm.impala"
        bool _30058;
        _30058 = 0 < shm_dimy_30057;
        #line 38 "gpu_device_shm.impala"
        if (_30058) goto l30059; else goto l30234;
    l30234: ;
        #line 231 "gpu_device_shm.impala"
        goto l30166;
    l30059: ;
        #line 226 "gpu_device_shm.impala"
        int _30102;
        _30102 = tidy_30042 * shm_dimx_30067;
        #line 221 "gpu_device_shm.impala"
        int _30086;
        _30086 = gid_y_30085 - extend_height_30055;
        #line 219 "gpu_device_shm.impala"
        int _30077;
        _30077 = gid_x_30076 - h_anchor_30065;
        #line 217 "gpu_device_shm.impala"
        bool _30073;
        _30073 = tidy_30042 < shm_dimy_30057;
        #line 224 "gpu_device_shm.impala"
        bool _30087;
        _30087 = 0 <= _30086;
        #line 224 "gpu_device_shm.impala"
        bool _30091;
        _30091 = _30086 < _30090;
        #line 227 "gpu_device_shm.impala"
        int _30096;
        _30096 = _30086 * _30081;
        #line 38 "gpu_device_shm.impala"
        p_30061 = 0;
        goto l30060;
    l30060: ;
        _30061 = p_30061;
        #line 38 "gpu_device_shm.impala"
        bool _30068;
        _30068 = _30061 < shm_dimx_30067;
        #line 38 "gpu_device_shm.impala"
        if (_30068) goto l30069; else goto l30115;
    l30115: ;
        #line 38 "gpu_device_shm.impala"
        p_30117 = bdimy_30018;
        goto l30116;
    l30116: ;
        _30117 = p_30117;
        #line 38 "gpu_device_shm.impala"
        bool _30119;
        _30119 = _30117 < shm_dimy_30057;
        #line 38 "gpu_device_shm.impala"
        if (_30119) goto l30120; else goto l30165;
    l30165: ;
        #line 231 "gpu_device_shm.impala"
        goto l30166;
    l30166: ;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        __syncthreads();
    l30171: ;
        #line 237 "gpu_device_shm.impala"
        bool _30173;
        _30173 = gid_x_30076 < _30081;
        #line 237 "gpu_device_shm.impala"
        if (_30173) goto l30174; else goto l30233;
    l30233: ;
        #line 240 "gpu_device_shm.impala"
        goto l30232;
    l30174: ;
        #line 237 "gpu_device_shm.impala"
        bool _30175;
        _30175 = gid_y_30085 < _30090;
        #line 237 "gpu_device_shm.impala"
        if (_30175) goto l30176; else goto l30231;
    l30231: ;
        #line 240 "gpu_device_shm.impala"
        goto l30232;
    l30232: ;
        return ;
    l30176: ;
        #line 64 "gpu_device_shm.impala"
        char* _30213;
        _30213 = _25296_30004.e1;
        #line 204 "gpu_device_shm.impala"
        int _30198;
        _30198 = extend_height_30055 - _30084;
        #line 64 "gpu_device_shm.impala"
        double* _30214;
        union { double* dst; char* src; } u_30214;
        u_30214.src = _30213;
        _30214 = u_30214.dst;
        #line 203 "gpu_device_shm.impala"
        int _30203;
        _30203 = h_anchor_30065 - _30075;
        #line 64 "gpu_device_shm.impala"
        int _30215;
        _30215 = gid_y_30085 * _30081;
        #line 17 "gaussian.impala"
        bool _30177;
        _30177 = h_anchor_30065 <= gid_x_30076;
        #line 64 "gpu_device_shm.impala"
        int _30216;
        _30216 = _30215 + gid_x_30076;
        #line 72 "gpu_device_shm.impala"
        int _30199;
        _30199 = gid_y_30085 + _30198;
        #line 64 "gpu_device_shm.impala"
        double* _30217;
        _30217 = _30214 + _30216;
        #line 72 "gpu_device_shm.impala"
        int _30200;
        _30200 = _30199 * shm_dimx_30067;
        #line 17 "gaussian.impala"
        if (_30177) goto l30178; else goto l30230;
    l30230: ;
        #line 27 "gaussian.impala"
        goto l30222;
    l30178: ;
        #line 17 "gaussian.impala"
        int _30179;
        _30179 = _30081 - h_anchor_30065;
        #line 17 "gaussian.impala"
        bool _30180;
        _30180 = gid_x_30076 < _30179;
        #line 17 "gaussian.impala"
        if (_30180) goto l30181; else goto l30221;
    l30221: ;
        #line 27 "gaussian.impala"
        goto l30222;
    l30222: ;
        #line 72 "gpu_device_shm.impala"
        int _30223;
        _30223 = _30200 + gid_x_30076;
        #line 72 "gpu_device_shm.impala"
        int _30224;
        _30224 = _30223 + _30203;
        #line 72 "gpu_device_shm.impala"
        double* _30225;
        _30225 = reserve_shared_30050 + _30224;
        #line 72 "gpu_device_shm.impala"
        double _30226;
        _30226 = *_30225;
        #line 72 "gpu_device_shm.impala"
        double _30228;
        _30228 = _30226;
        #line 64 "gpu_device_shm.impala"
        *_30217 = _30228;
        return ;
    l30181: ;
        #line 77 "gpu_device_shm.impala"
        struct_Buffer_6327 _30191;
        _30191 = _25298_30006.e0;
        #line 19 "gaussian.impala"
        int _30219;
        _30219 = 0 - h_anchor_30065;
        #line 77 "gpu_device_shm.impala"
        char* _30192;
        _30192 = _30191.e1;
        #line 19 "gaussian.impala"
        int _30187;
        _30187 = 1 + h_anchor_30065;
        #line 77 "gpu_device_shm.impala"
        double* _30193;
        union { double* dst; char* src; } u_30193;
        u_30193.src = _30192;
        _30193 = u_30193.dst;
        #line 27 "gpu_device_shm.impala"
        p_30183 = _30219;
        psum_30185 = 0.000000e+00;
        goto l30182;
    l30182: ;
        _30183 = p_30183;
        sum_30185 = psum_30185;
        #line 27 "gpu_device_shm.impala"
        bool _30188;
        _30188 = _30183 < _30187;
        #line 27 "gpu_device_shm.impala"
        if (_30188) goto l30189; else goto l30212;
    l30212: ;
        #line 64 "gpu_device_shm.impala"
        *_30217 = sum_30185;
        return ;
    l30189: ;
        #line 31 "gpu_device_shm.impala"
        int _30190;
        _30190 = 1 + _30183;
        #line 21 "gaussian.impala"
        int _30201;
        _30201 = gid_x_30076 + _30183;
        #line 72 "gpu_device_shm.impala"
        int _30202;
        _30202 = _30200 + _30201;
        #line 21 "gaussian.impala"
        int _30194;
        _30194 = _30183 + h_anchor_30065;
        #line 72 "gpu_device_shm.impala"
        int _30204;
        _30204 = _30202 + _30203;
        #line 76 "gpu_device_shm.impala"
        double* i_30195;
        i_30195 = _30193 + _30194;
        #line 72 "gpu_device_shm.impala"
        double* _30205;
        _30205 = reserve_shared_30050 + _30204;
        #line 77 "gpu_device_shm.impala"
        double _30196;
        _30196 = *i_30195;
        #line 77 "gpu_device_shm.impala"
        double _30208;
        _30208 = _30196;
        #line 72 "gpu_device_shm.impala"
        double _30206;
        _30206 = *_30205;
        #line 72 "gpu_device_shm.impala"
        double _30209;
        _30209 = _30206;
        #line 21 "gaussian.impala"
        double _30210;
        _30210 = _30208 * _30209;
        #line 21 "gaussian.impala"
        double _30211;
        _30211 = sum_30185 + _30210;
        #line 27 "gpu_device_shm.impala"
        p_30183 = _30190;
        psum_30185 = _30211;
        goto l30182;
    l30120: ;
        #line 221 "gpu_device_shm.impala"
        int img_index_y_30137;
        img_index_y_30137 = _30086 + _30117;
        #line 214 "gpu_device_shm.impala"
        int shm_index_y_30129;
        shm_index_y_30129 = tidy_30042 + _30117;
        #line 224 "gpu_device_shm.impala"
        bool _30138;
        _30138 = 0 <= img_index_y_30137;
        #line 227 "gpu_device_shm.impala"
        int _30145;
        _30145 = img_index_y_30137 * _30081;
        #line 224 "gpu_device_shm.impala"
        bool _30140;
        _30140 = img_index_y_30137 < _30090;
        #line 217 "gpu_device_shm.impala"
        bool _30130;
        _30130 = shm_index_y_30129 < shm_dimy_30057;
        #line 226 "gpu_device_shm.impala"
        int _30150;
        _30150 = shm_index_y_30129 * shm_dimx_30067;
        #line 38 "gpu_device_shm.impala"
        p_30122 = 0;
        goto l30121;
    l30121: ;
        _30122 = p_30122;
        #line 38 "gpu_device_shm.impala"
        bool _30124;
        _30124 = _30122 < shm_dimx_30067;
        #line 38 "gpu_device_shm.impala"
        if (_30124) goto l30125; else goto l30163;
    l30163: ;
        #line 42 "gpu_device_shm.impala"
        int _30164;
        _30164 = _30117 + bdimy_30018;
        #line 38 "gpu_device_shm.impala"
        p_30117 = _30164;
        goto l30116;
    l30125: ;
        #line 212 "gpu_device_shm.impala"
        int shm_index_x_30126;
        shm_index_x_30126 = tidx_30036 + _30122;
        #line 217 "gpu_device_shm.impala"
        bool _30127;
        _30127 = shm_index_x_30126 < shm_dimx_30067;
        #line 217 "gpu_device_shm.impala"
        if (_30127) goto l30128; else goto l30162;
    l30162: ;
        #line 229 "gpu_device_shm.impala"
        goto l30161;
    l30128: ;
        #line 217 "gpu_device_shm.impala"
        if (_30130) goto l30131; else goto l30160;
    l30160: ;
        #line 229 "gpu_device_shm.impala"
        goto l30161;
    l30161: ;
        #line 40 "gpu_device_shm.impala"
        goto l30142;
    l30131: ;
        #line 219 "gpu_device_shm.impala"
        int img_index_x_30132;
        img_index_x_30132 = _30077 + _30122;
        #line 224 "gpu_device_shm.impala"
        bool _30133;
        _30133 = 0 <= img_index_x_30132;
        #line 224 "gpu_device_shm.impala"
        if (_30133) goto l30134; else goto l30159;
    l30159: ;
        #line 228 "gpu_device_shm.impala"
        goto l30156;
    l30134: ;
        #line 224 "gpu_device_shm.impala"
        bool _30135;
        _30135 = img_index_x_30132 < _30081;
        #line 224 "gpu_device_shm.impala"
        if (_30135) goto l30136; else goto l30158;
    l30158: ;
        #line 228 "gpu_device_shm.impala"
        goto l30156;
    l30136: ;
        #line 224 "gpu_device_shm.impala"
        if (_30138) goto l30139; else goto l30157;
    l30157: ;
        #line 228 "gpu_device_shm.impala"
        goto l30156;
    l30139: ;
        #line 224 "gpu_device_shm.impala"
        if (_30140) goto l30141; else goto l30155;
    l30155: ;
        #line 228 "gpu_device_shm.impala"
        goto l30156;
    l30156: ;
        #line 40 "gpu_device_shm.impala"
        goto l30142;
    l30141: ;
        #line 226 "gpu_device_shm.impala"
        int _30151;
        _30151 = _30150 + shm_index_x_30126;
        #line 226 "gpu_device_shm.impala"
        double* _30152;
        _30152 = reserve_shared_30050 + _30151;
        #line 227 "gpu_device_shm.impala"
        int _30146;
        _30146 = _30145 + img_index_x_30132;
        #line 227 "gpu_device_shm.impala"
        double* _30147;
        _30147 = _25297_30005 + _30146;
        #line 227 "gpu_device_shm.impala"
        double _30148;
        _30148 = *_30147;
        #line 227 "gpu_device_shm.impala"
        double _30153;
        _30153 = _30148;
        #line 226 "gpu_device_shm.impala"
        *_30152 = _30153;
        #line 40 "gpu_device_shm.impala"
        goto l30142;
    l30142: ;
        #line 42 "gpu_device_shm.impala"
        int _30144;
        _30144 = _30122 + bdimx_30012;
        #line 38 "gpu_device_shm.impala"
        p_30122 = _30144;
        goto l30121;
    l30069: ;
        #line 212 "gpu_device_shm.impala"
        int shm_index_x_30070;
        shm_index_x_30070 = tidx_30036 + _30061;
        #line 217 "gpu_device_shm.impala"
        bool _30071;
        _30071 = shm_index_x_30070 < shm_dimx_30067;
        #line 217 "gpu_device_shm.impala"
        if (_30071) goto l30072; else goto l30114;
    l30114: ;
        #line 229 "gpu_device_shm.impala"
        goto l30113;
    l30072: ;
        #line 217 "gpu_device_shm.impala"
        if (_30073) goto l30074; else goto l30112;
    l30112: ;
        #line 229 "gpu_device_shm.impala"
        goto l30113;
    l30113: ;
        #line 40 "gpu_device_shm.impala"
        goto l30093;
    l30074: ;
        #line 219 "gpu_device_shm.impala"
        int img_index_x_30078;
        img_index_x_30078 = _30077 + _30061;
        #line 224 "gpu_device_shm.impala"
        bool _30079;
        _30079 = 0 <= img_index_x_30078;
        #line 224 "gpu_device_shm.impala"
        if (_30079) goto l30080; else goto l30111;
    l30111: ;
        #line 228 "gpu_device_shm.impala"
        goto l30108;
    l30080: ;
        #line 224 "gpu_device_shm.impala"
        bool _30082;
        _30082 = img_index_x_30078 < _30081;
        #line 224 "gpu_device_shm.impala"
        if (_30082) goto l30083; else goto l30110;
    l30110: ;
        #line 228 "gpu_device_shm.impala"
        goto l30108;
    l30083: ;
        #line 224 "gpu_device_shm.impala"
        if (_30087) goto l30088; else goto l30109;
    l30109: ;
        #line 228 "gpu_device_shm.impala"
        goto l30108;
    l30088: ;
        #line 224 "gpu_device_shm.impala"
        if (_30091) goto l30092; else goto l30107;
    l30107: ;
        #line 228 "gpu_device_shm.impala"
        goto l30108;
    l30108: ;
        #line 40 "gpu_device_shm.impala"
        goto l30093;
    l30092: ;
        #line 226 "gpu_device_shm.impala"
        int _30103;
        _30103 = _30102 + shm_index_x_30070;
        #line 226 "gpu_device_shm.impala"
        double* _30104;
        _30104 = reserve_shared_30050 + _30103;
        #line 227 "gpu_device_shm.impala"
        int _30097;
        _30097 = _30096 + img_index_x_30078;
        #line 227 "gpu_device_shm.impala"
        double* _30098;
        _30098 = _25297_30005 + _30097;
        #line 227 "gpu_device_shm.impala"
        double _30099;
        _30099 = *_30098;
        #line 227 "gpu_device_shm.impala"
        double _30105;
        _30105 = _30099;
        #line 226 "gpu_device_shm.impala"
        *_30104 = _30105;
        #line 40 "gpu_device_shm.impala"
        goto l30093;
    l30093: ;
        #line 42 "gpu_device_shm.impala"
        int _30095;
        _30095 = _30061 + bdimx_30012;
        #line 38 "gpu_device_shm.impala"
        p_30061 = _30095;
        goto l30060;
}

__global__ __launch_bounds__ (128 * 1 * 1) void lambda_25546(struct_image_6326 _25549_30238, struct_filter_6332 _25550_30239, double* _25551_30240, double* _25552_30241, struct_Buffer_6327 _25553_30242) {
    __shared__ double ds_img[134][7];
    int  bdimx_30245;
    int pbdimx_30245;
    int  bdimy_30248;
    int pbdimy_30248;
    int  bidx_30251;
    int pbidx_30251;
    int  bidy_30254;
    int pbidy_30254;
    int  tidx_30257;
    int ptidx_30257;
    int  tidy_30260;
    int ptidy_30260;
    double*  reserve_shared_30263;
    double* preserve_shared_30263;
    int  _30271;
    int p_30271;
    int  _30324;
    int p_30324;
    int  _30387;
    int p_30387;
    double  sum_30389;
    double psum_30389;
    int  _30329;
    int p_30329;
    #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
    bdimx_30245 = blockDim_x();
    pbdimx_30245 = bdimx_30245;
    l30243: ;
        bdimx_30245 = pbdimx_30245;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        bdimy_30248 = blockDim_y();
        pbdimy_30248 = bdimy_30248;
    l30246: ;
        bdimy_30248 = pbdimy_30248;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        bidx_30251 = blockIdx_x();
        pbidx_30251 = bidx_30251;
    l30249: ;
        bidx_30251 = pbidx_30251;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        bidy_30254 = blockIdx_y();
        pbidy_30254 = bidy_30254;
    l30252: ;
        bidy_30254 = pbidy_30254;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        tidx_30257 = threadIdx_x();
        ptidx_30257 = tidx_30257;
    l30255: ;
        tidx_30257 = ptidx_30257;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        tidy_30260 = threadIdx_y();
        ptidy_30260 = tidy_30260;
    l30258: ;
        tidy_30260 = ptidy_30260;
        #line 201 "gpu_device_shm.impala"
        __shared__ double reserver_reserve_shared_30263[938];
        preserve_shared_30263 = reserver_reserve_shared_30263;
    l30261: ;
        reserve_shared_30263 = preserve_shared_30263;
        #line 187 "gpu_device_shm.impala"
        int _30284;
        _30284 = bidx_30251 * bdimx_30245;
        #line 11 "main.impala"
        int _30298;
        _30298 = _25549_30238.e3;
        #line 187 "gpu_device_shm.impala"
        int gid_x_30285;
        gid_x_30285 = _30284 + tidx_30257;
        #line 6 "gaussian.impala"
        int _30264;
        _30264 = _25550_30239.e2;
        #line 191 "gpu_device_shm.impala"
        int _30273;
        _30273 = _25550_30239.e1;
        #line 189 "gpu_device_shm.impala"
        int _30293;
        _30293 = bidy_30254 * bdimy_30248;
        #line 11 "main.impala"
        int _30290;
        _30290 = _25549_30238.e2;
        #line 189 "gpu_device_shm.impala"
        int gid_y_30294;
        gid_y_30294 = _30293 + tidy_30260;
        #line 191 "gpu_device_shm.impala"
        int extend_width_30274;
        extend_width_30274 = _30273 / 2;
        #line 6 "gaussian.impala"
        int v_anchor_30265;
        v_anchor_30265 = _30264 / 2;
        #line 195 "gpu_device_shm.impala"
        int _30275;
        _30275 = 2 * extend_width_30274;
        #line 197 "gpu_device_shm.impala"
        int _30266;
        _30266 = 2 * v_anchor_30265;
        #line 195 "gpu_device_shm.impala"
        int shm_dimx_30276;
        shm_dimx_30276 = bdimx_30245 + _30275;
        #line 197 "gpu_device_shm.impala"
        int shm_dimy_30267;
        shm_dimy_30267 = bdimy_30248 + _30266;
        #line 38 "gpu_device_shm.impala"
        bool _30268;
        _30268 = 0 < shm_dimy_30267;
        #line 38 "gpu_device_shm.impala"
        if (_30268) goto l30269; else goto l30438;
    l30438: ;
        #line 231 "gpu_device_shm.impala"
        goto l30373;
    l30269: ;
        #line 226 "gpu_device_shm.impala"
        int _30309;
        _30309 = tidy_30260 * shm_dimx_30276;
        #line 219 "gpu_device_shm.impala"
        int _30286;
        _30286 = gid_x_30285 - extend_width_30274;
        #line 221 "gpu_device_shm.impala"
        int _30295;
        _30295 = gid_y_30294 - v_anchor_30265;
        #line 217 "gpu_device_shm.impala"
        bool _30282;
        _30282 = tidy_30260 < shm_dimy_30267;
        #line 224 "gpu_device_shm.impala"
        bool _30296;
        _30296 = 0 <= _30295;
        #line 224 "gpu_device_shm.impala"
        bool _30299;
        _30299 = _30295 < _30298;
        #line 227 "gpu_device_shm.impala"
        int _30304;
        _30304 = _30295 * _30290;
        #line 38 "gpu_device_shm.impala"
        p_30271 = 0;
        goto l30270;
    l30270: ;
        _30271 = p_30271;
        #line 38 "gpu_device_shm.impala"
        bool _30277;
        _30277 = _30271 < shm_dimx_30276;
        #line 38 "gpu_device_shm.impala"
        if (_30277) goto l30278; else goto l30322;
    l30322: ;
        #line 38 "gpu_device_shm.impala"
        p_30324 = bdimy_30248;
        goto l30323;
    l30323: ;
        _30324 = p_30324;
        #line 38 "gpu_device_shm.impala"
        bool _30326;
        _30326 = _30324 < shm_dimy_30267;
        #line 38 "gpu_device_shm.impala"
        if (_30326) goto l30327; else goto l30372;
    l30372: ;
        #line 231 "gpu_device_shm.impala"
        goto l30373;
    l30373: ;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        __syncthreads();
    l30375: ;
        #line 237 "gpu_device_shm.impala"
        bool _30377;
        _30377 = gid_x_30285 < _30290;
        #line 237 "gpu_device_shm.impala"
        if (_30377) goto l30378; else goto l30437;
    l30437: ;
        #line 240 "gpu_device_shm.impala"
        goto l30436;
    l30378: ;
        #line 237 "gpu_device_shm.impala"
        bool _30379;
        _30379 = gid_y_30294 < _30298;
        #line 237 "gpu_device_shm.impala"
        if (_30379) goto l30380; else goto l30435;
    l30435: ;
        #line 240 "gpu_device_shm.impala"
        goto l30436;
    l30436: ;
        return ;
    l30380: ;
        #line 204 "gpu_device_shm.impala"
        int _30402;
        _30402 = v_anchor_30265 - _30293;
        #line 64 "gpu_device_shm.impala"
        char* _30416;
        _30416 = _25553_30242.e1;
        #line 203 "gpu_device_shm.impala"
        int _30406;
        _30406 = extend_width_30274 - _30284;
        #line 64 "gpu_device_shm.impala"
        int _30418;
        _30418 = gid_y_30294 * _30290;
        #line 64 "gpu_device_shm.impala"
        int _30419;
        _30419 = _30418 + gid_x_30285;
        #line 64 "gpu_device_shm.impala"
        double* _30417;
        union { double* dst; char* src; } u_30417;
        u_30417.src = _30416;
        _30417 = u_30417.dst;
        #line 39 "gaussian.impala"
        bool _30381;
        _30381 = v_anchor_30265 <= gid_y_30294;
        #line 64 "gpu_device_shm.impala"
        double* _30420;
        _30420 = _30417 + _30419;
        #line 39 "gaussian.impala"
        if (_30381) goto l30382; else goto l30434;
    l30434: ;
        #line 49 "gaussian.impala"
        goto l30424;
    l30382: ;
        #line 39 "gaussian.impala"
        int _30383;
        _30383 = _30298 - v_anchor_30265;
        #line 39 "gaussian.impala"
        bool _30384;
        _30384 = gid_y_30294 < _30383;
        #line 39 "gaussian.impala"
        if (_30384) goto l30385; else goto l30423;
    l30423: ;
        #line 49 "gaussian.impala"
        goto l30424;
    l30424: ;
        #line 72 "gpu_device_shm.impala"
        int _30425;
        _30425 = gid_y_30294 + _30402;
        #line 72 "gpu_device_shm.impala"
        int _30426;
        _30426 = _30425 * shm_dimx_30276;
        #line 72 "gpu_device_shm.impala"
        int _30427;
        _30427 = _30426 + gid_x_30285;
        #line 72 "gpu_device_shm.impala"
        int _30428;
        _30428 = _30427 + _30406;
        #line 72 "gpu_device_shm.impala"
        double* _30429;
        _30429 = reserve_shared_30263 + _30428;
        #line 72 "gpu_device_shm.impala"
        double _30430;
        _30430 = *_30429;
        #line 72 "gpu_device_shm.impala"
        double _30432;
        _30432 = _30430;
        #line 64 "gpu_device_shm.impala"
        *_30420 = _30432;
        return ;
    l30385: ;
        #line 41 "gaussian.impala"
        int _30422;
        _30422 = 0 - v_anchor_30265;
        #line 77 "gpu_device_shm.impala"
        struct_Buffer_6327 _30394;
        _30394 = _25550_30239.e0;
        #line 41 "gaussian.impala"
        int _30390;
        _30390 = 1 + v_anchor_30265;
        #line 77 "gpu_device_shm.impala"
        char* _30395;
        _30395 = _30394.e1;
        #line 77 "gpu_device_shm.impala"
        double* _30396;
        union { double* dst; char* src; } u_30396;
        u_30396.src = _30395;
        _30396 = u_30396.dst;
        #line 27 "gpu_device_shm.impala"
        p_30387 = _30422;
        psum_30389 = 0.000000e+00;
        goto l30386;
    l30386: ;
        _30387 = p_30387;
        sum_30389 = psum_30389;
        #line 27 "gpu_device_shm.impala"
        bool _30391;
        _30391 = _30387 < _30390;
        #line 27 "gpu_device_shm.impala"
        if (_30391) goto l30392; else goto l30415;
    l30415: ;
        #line 64 "gpu_device_shm.impala"
        *_30420 = sum_30389;
        return ;
    l30392: ;
        #line 43 "gaussian.impala"
        int _30401;
        _30401 = gid_y_30294 + _30387;
        #line 43 "gaussian.impala"
        int _30397;
        _30397 = _30387 + v_anchor_30265;
        #line 31 "gpu_device_shm.impala"
        int _30393;
        _30393 = 1 + _30387;
        #line 72 "gpu_device_shm.impala"
        int _30403;
        _30403 = _30401 + _30402;
        #line 76 "gpu_device_shm.impala"
        double* i_30398;
        i_30398 = _30396 + _30397;
        #line 72 "gpu_device_shm.impala"
        int _30404;
        _30404 = _30403 * shm_dimx_30276;
        #line 77 "gpu_device_shm.impala"
        double _30399;
        _30399 = *i_30398;
        #line 72 "gpu_device_shm.impala"
        int _30405;
        _30405 = _30404 + gid_x_30285;
        #line 77 "gpu_device_shm.impala"
        double _30411;
        _30411 = _30399;
        #line 72 "gpu_device_shm.impala"
        int _30407;
        _30407 = _30405 + _30406;
        #line 72 "gpu_device_shm.impala"
        double* _30408;
        _30408 = reserve_shared_30263 + _30407;
        #line 72 "gpu_device_shm.impala"
        double _30409;
        _30409 = *_30408;
        #line 72 "gpu_device_shm.impala"
        double _30412;
        _30412 = _30409;
        #line 43 "gaussian.impala"
        double _30413;
        _30413 = _30411 * _30412;
        #line 43 "gaussian.impala"
        double _30414;
        _30414 = sum_30389 + _30413;
        #line 27 "gpu_device_shm.impala"
        p_30387 = _30393;
        psum_30389 = _30414;
        goto l30386;
    l30327: ;
        #line 221 "gpu_device_shm.impala"
        int img_index_y_30344;
        img_index_y_30344 = _30295 + _30324;
        #line 224 "gpu_device_shm.impala"
        bool _30347;
        _30347 = img_index_y_30344 < _30298;
        #line 224 "gpu_device_shm.impala"
        bool _30345;
        _30345 = 0 <= img_index_y_30344;
        #line 214 "gpu_device_shm.impala"
        int shm_index_y_30336;
        shm_index_y_30336 = tidy_30260 + _30324;
        #line 227 "gpu_device_shm.impala"
        int _30352;
        _30352 = img_index_y_30344 * _30290;
        #line 217 "gpu_device_shm.impala"
        bool _30337;
        _30337 = shm_index_y_30336 < shm_dimy_30267;
        #line 226 "gpu_device_shm.impala"
        int _30357;
        _30357 = shm_index_y_30336 * shm_dimx_30276;
        #line 38 "gpu_device_shm.impala"
        p_30329 = 0;
        goto l30328;
    l30328: ;
        _30329 = p_30329;
        #line 38 "gpu_device_shm.impala"
        bool _30331;
        _30331 = _30329 < shm_dimx_30276;
        #line 38 "gpu_device_shm.impala"
        if (_30331) goto l30332; else goto l30370;
    l30370: ;
        #line 42 "gpu_device_shm.impala"
        int _30371;
        _30371 = _30324 + bdimy_30248;
        #line 38 "gpu_device_shm.impala"
        p_30324 = _30371;
        goto l30323;
    l30332: ;
        #line 212 "gpu_device_shm.impala"
        int shm_index_x_30333;
        shm_index_x_30333 = tidx_30257 + _30329;
        #line 217 "gpu_device_shm.impala"
        bool _30334;
        _30334 = shm_index_x_30333 < shm_dimx_30276;
        #line 217 "gpu_device_shm.impala"
        if (_30334) goto l30335; else goto l30369;
    l30369: ;
        #line 229 "gpu_device_shm.impala"
        goto l30368;
    l30335: ;
        #line 217 "gpu_device_shm.impala"
        if (_30337) goto l30338; else goto l30367;
    l30367: ;
        #line 229 "gpu_device_shm.impala"
        goto l30368;
    l30368: ;
        #line 40 "gpu_device_shm.impala"
        goto l30349;
    l30338: ;
        #line 219 "gpu_device_shm.impala"
        int img_index_x_30339;
        img_index_x_30339 = _30286 + _30329;
        #line 224 "gpu_device_shm.impala"
        bool _30340;
        _30340 = 0 <= img_index_x_30339;
        #line 224 "gpu_device_shm.impala"
        if (_30340) goto l30341; else goto l30366;
    l30366: ;
        #line 228 "gpu_device_shm.impala"
        goto l30363;
    l30341: ;
        #line 224 "gpu_device_shm.impala"
        bool _30342;
        _30342 = img_index_x_30339 < _30290;
        #line 224 "gpu_device_shm.impala"
        if (_30342) goto l30343; else goto l30365;
    l30365: ;
        #line 228 "gpu_device_shm.impala"
        goto l30363;
    l30343: ;
        #line 224 "gpu_device_shm.impala"
        if (_30345) goto l30346; else goto l30364;
    l30364: ;
        #line 228 "gpu_device_shm.impala"
        goto l30363;
    l30346: ;
        #line 224 "gpu_device_shm.impala"
        if (_30347) goto l30348; else goto l30362;
    l30362: ;
        #line 228 "gpu_device_shm.impala"
        goto l30363;
    l30363: ;
        #line 40 "gpu_device_shm.impala"
        goto l30349;
    l30348: ;
        #line 227 "gpu_device_shm.impala"
        int _30353;
        _30353 = _30352 + img_index_x_30339;
        #line 227 "gpu_device_shm.impala"
        double* _30354;
        _30354 = _25551_30240 + _30353;
        #line 227 "gpu_device_shm.impala"
        double _30355;
        _30355 = *_30354;
        #line 226 "gpu_device_shm.impala"
        int _30358;
        _30358 = _30357 + shm_index_x_30333;
        #line 227 "gpu_device_shm.impala"
        double _30360;
        _30360 = _30355;
        #line 226 "gpu_device_shm.impala"
        double* _30359;
        _30359 = reserve_shared_30263 + _30358;
        #line 226 "gpu_device_shm.impala"
        *_30359 = _30360;
        #line 40 "gpu_device_shm.impala"
        goto l30349;
    l30349: ;
        #line 42 "gpu_device_shm.impala"
        int _30351;
        _30351 = _30329 + bdimx_30245;
        #line 38 "gpu_device_shm.impala"
        p_30329 = _30351;
        goto l30328;
    l30278: ;
        #line 212 "gpu_device_shm.impala"
        int shm_index_x_30279;
        shm_index_x_30279 = tidx_30257 + _30271;
        #line 217 "gpu_device_shm.impala"
        bool _30280;
        _30280 = shm_index_x_30279 < shm_dimx_30276;
        #line 217 "gpu_device_shm.impala"
        if (_30280) goto l30281; else goto l30321;
    l30321: ;
        #line 229 "gpu_device_shm.impala"
        goto l30320;
    l30281: ;
        #line 217 "gpu_device_shm.impala"
        if (_30282) goto l30283; else goto l30319;
    l30319: ;
        #line 229 "gpu_device_shm.impala"
        goto l30320;
    l30320: ;
        #line 40 "gpu_device_shm.impala"
        goto l30301;
    l30283: ;
        #line 219 "gpu_device_shm.impala"
        int img_index_x_30287;
        img_index_x_30287 = _30286 + _30271;
        #line 224 "gpu_device_shm.impala"
        bool _30288;
        _30288 = 0 <= img_index_x_30287;
        #line 224 "gpu_device_shm.impala"
        if (_30288) goto l30289; else goto l30318;
    l30318: ;
        #line 228 "gpu_device_shm.impala"
        goto l30315;
    l30289: ;
        #line 224 "gpu_device_shm.impala"
        bool _30291;
        _30291 = img_index_x_30287 < _30290;
        #line 224 "gpu_device_shm.impala"
        if (_30291) goto l30292; else goto l30317;
    l30317: ;
        #line 228 "gpu_device_shm.impala"
        goto l30315;
    l30292: ;
        #line 224 "gpu_device_shm.impala"
        if (_30296) goto l30297; else goto l30316;
    l30316: ;
        #line 228 "gpu_device_shm.impala"
        goto l30315;
    l30297: ;
        #line 224 "gpu_device_shm.impala"
        if (_30299) goto l30300; else goto l30314;
    l30314: ;
        #line 228 "gpu_device_shm.impala"
        goto l30315;
    l30315: ;
        #line 40 "gpu_device_shm.impala"
        goto l30301;
    l30300: ;
        #line 226 "gpu_device_shm.impala"
        int _30310;
        _30310 = _30309 + shm_index_x_30279;
        #line 227 "gpu_device_shm.impala"
        int _30305;
        _30305 = _30304 + img_index_x_30287;
        #line 226 "gpu_device_shm.impala"
        double* _30311;
        _30311 = reserve_shared_30263 + _30310;
        #line 227 "gpu_device_shm.impala"
        double* _30306;
        _30306 = _25551_30240 + _30305;
        #line 227 "gpu_device_shm.impala"
        double _30307;
        _30307 = *_30306;
        #line 227 "gpu_device_shm.impala"
        double _30312;
        _30312 = _30307;
        #line 226 "gpu_device_shm.impala"
        *_30311 = _30312;
        #line 40 "gpu_device_shm.impala"
        goto l30301;
    l30301: ;
        #line 42 "gpu_device_shm.impala"
        int _30303;
        _30303 = _30271 + bdimx_30245;
        #line 38 "gpu_device_shm.impala"
        p_30271 = _30303;
        goto l30270;
}

}
	.file	"tmpxft_00228919_00000000-6_main_shm.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2022:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2022:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z91__device_stub__Z12lambda_2529217struct_image_632618struct_Buffer_6327Pd18struct_filter_6332R17struct_image_6326R18struct_Buffer_6327PdR18struct_filter_6332
	.type	_Z91__device_stub__Z12lambda_2529217struct_image_632618struct_Buffer_6327Pd18struct_filter_6332R17struct_image_6326R18struct_Buffer_6327PdR18struct_filter_6332, @function
_Z91__device_stub__Z12lambda_2529217struct_image_632618struct_Buffer_6327Pd18struct_filter_6332R17struct_image_6326R18struct_Buffer_6327PdR18struct_filter_6332:
.LFB2044:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdx, 8(%rsp)
	leaq	24(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movq	%rdi, 88(%rsp)
	leaq	40(%rsp), %rdi
	movq	%rsi, 96(%rsp)
	leaq	52(%rsp), %rsi
	movq	%rax, 104(%rsp)
	movabsq	$4294967297, %rax
	movq	%rcx, 112(%rsp)
	leaq	32(%rsp), %rcx
	movq	%rax, 40(%rsp)
	movl	$1, 48(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	32(%rsp)
	.cfi_def_cfa_offset 152
	leaq	lambda_25292(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 160
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L2:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2044:
	.size	_Z91__device_stub__Z12lambda_2529217struct_image_632618struct_Buffer_6327Pd18struct_filter_6332R17struct_image_6326R18struct_Buffer_6327PdR18struct_filter_6332, .-_Z91__device_stub__Z12lambda_2529217struct_image_632618struct_Buffer_6327Pd18struct_filter_6332R17struct_image_6326R18struct_Buffer_6327PdR18struct_filter_6332
	.globl	lambda_25292
	.type	lambda_25292, @function
lambda_25292:
.LFB2045:
	.cfi_startproc
	endbr64
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movq	%rdi, (%rsp)
	leaq	72(%rsp), %rcx
	leaq	32(%rsp), %rdi
	movq	%rsi, 8(%rsp)
	movq	%rsp, %rsi
	call	_Z91__device_stub__Z12lambda_2529217struct_image_632618struct_Buffer_6327Pd18struct_filter_6332R17struct_image_6326R18struct_Buffer_6327PdR18struct_filter_6332
	addq	$24, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2045:
	.size	lambda_25292, .-lambda_25292
	.globl	_Z94__device_stub__Z12lambda_2554617struct_image_632618struct_filter_6332PdS1_18struct_Buffer_6327R17struct_image_6326R18struct_filter_6332PdS3_R18struct_Buffer_6327
	.type	_Z94__device_stub__Z12lambda_2554617struct_image_632618struct_filter_6332PdS1_18struct_Buffer_6327R17struct_image_6326R18struct_filter_6332PdS3_R18struct_Buffer_6327, @function
_Z94__device_stub__Z12lambda_2554617struct_image_632618struct_filter_6332PdS1_18struct_Buffer_6327R17struct_image_6326R18struct_filter_6332PdS3_R18struct_Buffer_6327:
.LFB2046:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdx, 8(%rsp)
	leaq	16(%rsp), %rdx
	movq	%rcx, (%rsp)
	leaq	24(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movq	%rdi, 80(%rsp)
	leaq	32(%rsp), %rdi
	movq	%rax, 96(%rsp)
	movq	%rsp, %rax
	movq	%rsi, 88(%rsp)
	leaq	44(%rsp), %rsi
	movq	%rax, 104(%rsp)
	movabsq	$4294967297, %rax
	movq	%r8, 112(%rsp)
	movq	%rax, 32(%rsp)
	movl	$1, 40(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L9
	pushq	24(%rsp)
	.cfi_def_cfa_offset 152
	leaq	lambda_25546(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 160
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	96(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L9:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L11
	call	__stack_chk_fail@PLT
.L11:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2046:
	.size	_Z94__device_stub__Z12lambda_2554617struct_image_632618struct_filter_6332PdS1_18struct_Buffer_6327R17struct_image_6326R18struct_filter_6332PdS3_R18struct_Buffer_6327, .-_Z94__device_stub__Z12lambda_2554617struct_image_632618struct_filter_6332PdS1_18struct_Buffer_6327R17struct_image_6326R18struct_filter_6332PdS3_R18struct_Buffer_6327
	.globl	lambda_25546
	.type	lambda_25546, @function
lambda_25546:
.LFB2047:
	.cfi_startproc
	endbr64
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	leaq	72(%rsp), %rax
	movq	%rdx, (%rsp)
	movq	%rsp, %r8
	movq	%rdi, %rdx
	movq	%rcx, 8(%rsp)
	leaq	32(%rsp), %rdi
	movq	%rsi, %rcx
	movq	%rax, %rsi
	call	_Z94__device_stub__Z12lambda_2554617struct_image_632618struct_filter_6332PdS1_18struct_Buffer_6327R17struct_image_6326R18struct_filter_6332PdS3_R18struct_Buffer_6327
	addq	$24, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2047:
	.size	lambda_25546, .-lambda_25546
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"lambda_25546"
.LC1:
	.string	"lambda_25292"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2049:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC0(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	lambda_25546(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC1(%rip), %rdx
	orl	$-1, %r8d
	leaq	lambda_25292(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2049:
	.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>
extern "C" {
typedef struct {
    int e0;
    char* e1;
} struct_Buffer_6327;
typedef struct {
    struct_Buffer_6327 e0;
    struct_Buffer_6327 e1;
    int e2;
    int e3;
} struct_image_6326;
typedef struct {
    struct_Buffer_6327 e0;
    int e1;
    int e2;
} struct_filter_6332;

__device__ inline int threadIdx_x() { return threadIdx.x; }
__device__ inline int threadIdx_y() { return threadIdx.y; }
__device__ inline int threadIdx_z() { return threadIdx.z; }
__device__ inline int blockIdx_x() { return blockIdx.x; }
__device__ inline int blockIdx_y() { return blockIdx.y; }
__device__ inline int blockIdx_z() { return blockIdx.z; }
__device__ inline int blockDim_x() { return blockDim.x; }
__device__ inline int blockDim_y() { return blockDim.y; }
__device__ inline int blockDim_z() { return blockDim.z; }
__device__ inline int gridDim_x() { return gridDim.x; }
__device__ inline int gridDim_y() { return gridDim.y; }
__device__ inline int gridDim_z() { return gridDim.z; }
__global__ void lambda_25292(struct_image_6326, struct_Buffer_6327, double*, struct_filter_6332);
__global__ void lambda_25546(struct_image_6326, struct_filter_6332, double*, double*, struct_Buffer_6327);

__global__ __launch_bounds__ (128 * 1 * 1) void lambda_25292(struct_image_6326 _25295_30003, struct_Buffer_6327 _25296_30004, double* _25297_30005, struct_filter_6332 _25298_30006) {
    __shared__ double ds_img[134][7];
    int  bdimx_30012;
    int pbdimx_30012;
    int  bdimy_30018;
    int pbdimy_30018;
    int  bidx_30024;
    int pbidx_30024;
    int  bidy_30030;
    int pbidy_30030;
    int  tidx_30036;
    int ptidx_30036;
    int  tidy_30042;
    int ptidy_30042;
    double*  reserve_shared_30050;
    double* preserve_shared_30050;
    int  _30061;
    int p_30061;
    int  _30117;
    int p_30117;
    int  _30183;
    int p_30183;
    double  sum_30185;
    double psum_30185;
    int  _30122;
    int p_30122;
    #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
    bdimx_30012 = blockDim_x();
    pbdimx_30012 = bdimx_30012;
    l30010: ;
        bdimx_30012 = pbdimx_30012;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        bdimy_30018 = blockDim_y();
        pbdimy_30018 = bdimy_30018;
    l30016: ;
        bdimy_30018 = pbdimy_30018;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        bidx_30024 = blockIdx_x();
        pbidx_30024 = bidx_30024;
    l30022: ;
        bidx_30024 = pbidx_30024;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        bidy_30030 = blockIdx_y();
        pbidy_30030 = bidy_30030;
    l30028: ;
        bidy_30030 = pbidy_30030;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        tidx_30036 = threadIdx_x();
        ptidx_30036 = tidx_30036;
    l30034: ;
        tidx_30036 = ptidx_30036;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        tidy_30042 = threadIdx_y();
        ptidy_30042 = tidy_30042;
    l30040: ;
        tidy_30042 = ptidy_30042;
        #line 201 "gpu_device_shm.impala"
        __shared__ double reserver_reserve_shared_30050[938];
        preserve_shared_30050 = reserver_reserve_shared_30050;
    l30048: ;
        reserve_shared_30050 = preserve_shared_30050;
        #line 193 "gpu_device_shm.impala"
        int _30054;
        _30054 = _25298_30006.e2;
        #line 189 "gpu_device_shm.impala"
        int _30084;
        _30084 = bidy_30030 * bdimy_30018;
        #line 11 "main.impala"
        int _30081;
        _30081 = _25295_30003.e2;
        #line 187 "gpu_device_shm.impala"
        int _30075;
        _30075 = bidx_30024 * bdimx_30012;
        #line 11 "main.impala"
        int _30090;
        _30090 = _25295_30003.e3;
        #line 4 "gaussian.impala"
        int _30064;
        _30064 = _25298_30006.e1;
        #line 193 "gpu_device_shm.impala"
        int extend_height_30055;
        extend_height_30055 = _30054 / 2;
        #line 189 "gpu_device_shm.impala"
        int gid_y_30085;
        gid_y_30085 = _30084 + tidy_30042;
        #line 187 "gpu_device_shm.impala"
        int gid_x_30076;
        gid_x_30076 = _30075 + tidx_30036;
        #line 4 "gaussian.impala"
        int h_anchor_30065;
        h_anchor_30065 = _30064 / 2;
        #line 197 "gpu_device_shm.impala"
        int _30056;
        _30056 = 2 * extend_height_30055;
        #line 195 "gpu_device_shm.impala"
        int _30066;
        _30066 = 2 * h_anchor_30065;
        #line 197 "gpu_device_shm.impala"
        int shm_dimy_30057;
        shm_dimy_30057 = bdimy_30018 + _30056;
        #line 195 "gpu_device_shm.impala"
        int shm_dimx_30067;
        shm_dimx_30067 = bdimx_30012 + _30066;
        #line 38 "gpu_device_shm.impala"
        bool _30058;
        _30058 = 0 < shm_dimy_30057;
        #line 38 "gpu_device_shm.impala"
        if (_30058) goto l30059; else goto l30234;
    l30234: ;
        #line 231 "gpu_device_shm.impala"
        goto l30166;
    l30059: ;
        #line 226 "gpu_device_shm.impala"
        int _30102;
        _30102 = tidy_30042 * shm_dimx_30067;
        #line 221 "gpu_device_shm.impala"
        int _30086;
        _30086 = gid_y_30085 - extend_height_30055;
        #line 219 "gpu_device_shm.impala"
        int _30077;
        _30077 = gid_x_30076 - h_anchor_30065;
        #line 217 "gpu_device_shm.impala"
        bool _30073;
        _30073 = tidy_30042 < shm_dimy_30057;
        #line 224 "gpu_device_shm.impala"
        bool _30087;
        _30087 = 0 <= _30086;
        #line 224 "gpu_device_shm.impala"
        bool _30091;
        _30091 = _30086 < _30090;
        #line 227 "gpu_device_shm.impala"
        int _30096;
        _30096 = _30086 * _30081;
        #line 38 "gpu_device_shm.impala"
        p_30061 = 0;
        goto l30060;
    l30060: ;
        _30061 = p_30061;
        #line 38 "gpu_device_shm.impala"
        bool _30068;
        _30068 = _30061 < shm_dimx_30067;
        #line 38 "gpu_device_shm.impala"
        if (_30068) goto l30069; else goto l30115;
    l30115: ;
        #line 38 "gpu_device_shm.impala"
        p_30117 = bdimy_30018;
        goto l30116;
    l30116: ;
        _30117 = p_30117;
        #line 38 "gpu_device_shm.impala"
        bool _30119;
        _30119 = _30117 < shm_dimy_30057;
        #line 38 "gpu_device_shm.impala"
        if (_30119) goto l30120; else goto l30165;
    l30165: ;
        #line 231 "gpu_device_shm.impala"
        goto l30166;
    l30166: ;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        __syncthreads();
    l30171: ;
        #line 237 "gpu_device_shm.impala"
        bool _30173;
        _30173 = gid_x_30076 < _30081;
        #line 237 "gpu_device_shm.impala"
        if (_30173) goto l30174; else goto l30233;
    l30233: ;
        #line 240 "gpu_device_shm.impala"
        goto l30232;
    l30174: ;
        #line 237 "gpu_device_shm.impala"
        bool _30175;
        _30175 = gid_y_30085 < _30090;
        #line 237 "gpu_device_shm.impala"
        if (_30175) goto l30176; else goto l30231;
    l30231: ;
        #line 240 "gpu_device_shm.impala"
        goto l30232;
    l30232: ;
        return ;
    l30176: ;
        #line 64 "gpu_device_shm.impala"
        char* _30213;
        _30213 = _25296_30004.e1;
        #line 204 "gpu_device_shm.impala"
        int _30198;
        _30198 = extend_height_30055 - _30084;
        #line 64 "gpu_device_shm.impala"
        double* _30214;
        union { double* dst; char* src; } u_30214;
        u_30214.src = _30213;
        _30214 = u_30214.dst;
        #line 203 "gpu_device_shm.impala"
        int _30203;
        _30203 = h_anchor_30065 - _30075;
        #line 64 "gpu_device_shm.impala"
        int _30215;
        _30215 = gid_y_30085 * _30081;
        #line 17 "gaussian.impala"
        bool _30177;
        _30177 = h_anchor_30065 <= gid_x_30076;
        #line 64 "gpu_device_shm.impala"
        int _30216;
        _30216 = _30215 + gid_x_30076;
        #line 72 "gpu_device_shm.impala"
        int _30199;
        _30199 = gid_y_30085 + _30198;
        #line 64 "gpu_device_shm.impala"
        double* _30217;
        _30217 = _30214 + _30216;
        #line 72 "gpu_device_shm.impala"
        int _30200;
        _30200 = _30199 * shm_dimx_30067;
        #line 17 "gaussian.impala"
        if (_30177) goto l30178; else goto l30230;
    l30230: ;
        #line 27 "gaussian.impala"
        goto l30222;
    l30178: ;
        #line 17 "gaussian.impala"
        int _30179;
        _30179 = _30081 - h_anchor_30065;
        #line 17 "gaussian.impala"
        bool _30180;
        _30180 = gid_x_30076 < _30179;
        #line 17 "gaussian.impala"
        if (_30180) goto l30181; else goto l30221;
    l30221: ;
        #line 27 "gaussian.impala"
        goto l30222;
    l30222: ;
        #line 72 "gpu_device_shm.impala"
        int _30223;
        _30223 = _30200 + gid_x_30076;
        #line 72 "gpu_device_shm.impala"
        int _30224;
        _30224 = _30223 + _30203;
        #line 72 "gpu_device_shm.impala"
        double* _30225;
        _30225 = reserve_shared_30050 + _30224;
        #line 72 "gpu_device_shm.impala"
        double _30226;
        _30226 = *_30225;
        #line 72 "gpu_device_shm.impala"
        double _30228;
        _30228 = _30226;
        #line 64 "gpu_device_shm.impala"
        *_30217 = _30228;
        return ;
    l30181: ;
        #line 77 "gpu_device_shm.impala"
        struct_Buffer_6327 _30191;
        _30191 = _25298_30006.e0;
        #line 19 "gaussian.impala"
        int _30219;
        _30219 = 0 - h_anchor_30065;
        #line 77 "gpu_device_shm.impala"
        char* _30192;
        _30192 = _30191.e1;
        #line 19 "gaussian.impala"
        int _30187;
        _30187 = 1 + h_anchor_30065;
        #line 77 "gpu_device_shm.impala"
        double* _30193;
        union { double* dst; char* src; } u_30193;
        u_30193.src = _30192;
        _30193 = u_30193.dst;
        #line 27 "gpu_device_shm.impala"
        p_30183 = _30219;
        psum_30185 = 0.000000e+00;
        goto l30182;
    l30182: ;
        _30183 = p_30183;
        sum_30185 = psum_30185;
        #line 27 "gpu_device_shm.impala"
        bool _30188;
        _30188 = _30183 < _30187;
        #line 27 "gpu_device_shm.impala"
        if (_30188) goto l30189; else goto l30212;
    l30212: ;
        #line 64 "gpu_device_shm.impala"
        *_30217 = sum_30185;
        return ;
    l30189: ;
        #line 31 "gpu_device_shm.impala"
        int _30190;
        _30190 = 1 + _30183;
        #line 21 "gaussian.impala"
        int _30201;
        _30201 = gid_x_30076 + _30183;
        #line 72 "gpu_device_shm.impala"
        int _30202;
        _30202 = _30200 + _30201;
        #line 21 "gaussian.impala"
        int _30194;
        _30194 = _30183 + h_anchor_30065;
        #line 72 "gpu_device_shm.impala"
        int _30204;
        _30204 = _30202 + _30203;
        #line 76 "gpu_device_shm.impala"
        double* i_30195;
        i_30195 = _30193 + _30194;
        #line 72 "gpu_device_shm.impala"
        double* _30205;
        _30205 = reserve_shared_30050 + _30204;
        #line 77 "gpu_device_shm.impala"
        double _30196;
        _30196 = *i_30195;
        #line 77 "gpu_device_shm.impala"
        double _30208;
        _30208 = _30196;
        #line 72 "gpu_device_shm.impala"
        double _30206;
        _30206 = *_30205;
        #line 72 "gpu_device_shm.impala"
        double _30209;
        _30209 = _30206;
        #line 21 "gaussian.impala"
        double _30210;
        _30210 = _30208 * _30209;
        #line 21 "gaussian.impala"
        double _30211;
        _30211 = sum_30185 + _30210;
        #line 27 "gpu_device_shm.impala"
        p_30183 = _30190;
        psum_30185 = _30211;
        goto l30182;
    l30120: ;
        #line 221 "gpu_device_shm.impala"
        int img_index_y_30137;
        img_index_y_30137 = _30086 + _30117;
        #line 214 "gpu_device_shm.impala"
        int shm_index_y_30129;
        shm_index_y_30129 = tidy_30042 + _30117;
        #line 224 "gpu_device_shm.impala"
        bool _30138;
        _30138 = 0 <= img_index_y_30137;
        #line 227 "gpu_device_shm.impala"
        int _30145;
        _30145 = img_index_y_30137 * _30081;
        #line 224 "gpu_device_shm.impala"
        bool _30140;
        _30140 = img_index_y_30137 < _30090;
        #line 217 "gpu_device_shm.impala"
        bool _30130;
        _30130 = shm_index_y_30129 < shm_dimy_30057;
        #line 226 "gpu_device_shm.impala"
        int _30150;
        _30150 = shm_index_y_30129 * shm_dimx_30067;
        #line 38 "gpu_device_shm.impala"
        p_30122 = 0;
        goto l30121;
    l30121: ;
        _30122 = p_30122;
        #line 38 "gpu_device_shm.impala"
        bool _30124;
        _30124 = _30122 < shm_dimx_30067;
        #line 38 "gpu_device_shm.impala"
        if (_30124) goto l30125; else goto l30163;
    l30163: ;
        #line 42 "gpu_device_shm.impala"
        int _30164;
        _30164 = _30117 + bdimy_30018;
        #line 38 "gpu_device_shm.impala"
        p_30117 = _30164;
        goto l30116;
    l30125: ;
        #line 212 "gpu_device_shm.impala"
        int shm_index_x_30126;
        shm_index_x_30126 = tidx_30036 + _30122;
        #line 217 "gpu_device_shm.impala"
        bool _30127;
        _30127 = shm_index_x_30126 < shm_dimx_30067;
        #line 217 "gpu_device_shm.impala"
        if (_30127) goto l30128; else goto l30162;
    l30162: ;
        #line 229 "gpu_device_shm.impala"
        goto l30161;
    l30128: ;
        #line 217 "gpu_device_shm.impala"
        if (_30130) goto l30131; else goto l30160;
    l30160: ;
        #line 229 "gpu_device_shm.impala"
        goto l30161;
    l30161: ;
        #line 40 "gpu_device_shm.impala"
        goto l30142;
    l30131: ;
        #line 219 "gpu_device_shm.impala"
        int img_index_x_30132;
        img_index_x_30132 = _30077 + _30122;
        #line 224 "gpu_device_shm.impala"
        bool _30133;
        _30133 = 0 <= img_index_x_30132;
        #line 224 "gpu_device_shm.impala"
        if (_30133) goto l30134; else goto l30159;
    l30159: ;
        #line 228 "gpu_device_shm.impala"
        goto l30156;
    l30134: ;
        #line 224 "gpu_device_shm.impala"
        bool _30135;
        _30135 = img_index_x_30132 < _30081;
        #line 224 "gpu_device_shm.impala"
        if (_30135) goto l30136; else goto l30158;
    l30158: ;
        #line 228 "gpu_device_shm.impala"
        goto l30156;
    l30136: ;
        #line 224 "gpu_device_shm.impala"
        if (_30138) goto l30139; else goto l30157;
    l30157: ;
        #line 228 "gpu_device_shm.impala"
        goto l30156;
    l30139: ;
        #line 224 "gpu_device_shm.impala"
        if (_30140) goto l30141; else goto l30155;
    l30155: ;
        #line 228 "gpu_device_shm.impala"
        goto l30156;
    l30156: ;
        #line 40 "gpu_device_shm.impala"
        goto l30142;
    l30141: ;
        #line 226 "gpu_device_shm.impala"
        int _30151;
        _30151 = _30150 + shm_index_x_30126;
        #line 226 "gpu_device_shm.impala"
        double* _30152;
        _30152 = reserve_shared_30050 + _30151;
        #line 227 "gpu_device_shm.impala"
        int _30146;
        _30146 = _30145 + img_index_x_30132;
        #line 227 "gpu_device_shm.impala"
        double* _30147;
        _30147 = _25297_30005 + _30146;
        #line 227 "gpu_device_shm.impala"
        double _30148;
        _30148 = *_30147;
        #line 227 "gpu_device_shm.impala"
        double _30153;
        _30153 = _30148;
        #line 226 "gpu_device_shm.impala"
        *_30152 = _30153;
        #line 40 "gpu_device_shm.impala"
        goto l30142;
    l30142: ;
        #line 42 "gpu_device_shm.impala"
        int _30144;
        _30144 = _30122 + bdimx_30012;
        #line 38 "gpu_device_shm.impala"
        p_30122 = _30144;
        goto l30121;
    l30069: ;
        #line 212 "gpu_device_shm.impala"
        int shm_index_x_30070;
        shm_index_x_30070 = tidx_30036 + _30061;
        #line 217 "gpu_device_shm.impala"
        bool _30071;
        _30071 = shm_index_x_30070 < shm_dimx_30067;
        #line 217 "gpu_device_shm.impala"
        if (_30071) goto l30072; else goto l30114;
    l30114: ;
        #line 229 "gpu_device_shm.impala"
        goto l30113;
    l30072: ;
        #line 217 "gpu_device_shm.impala"
        if (_30073) goto l30074; else goto l30112;
    l30112: ;
        #line 229 "gpu_device_shm.impala"
        goto l30113;
    l30113: ;
        #line 40 "gpu_device_shm.impala"
        goto l30093;
    l30074: ;
        #line 219 "gpu_device_shm.impala"
        int img_index_x_30078;
        img_index_x_30078 = _30077 + _30061;
        #line 224 "gpu_device_shm.impala"
        bool _30079;
        _30079 = 0 <= img_index_x_30078;
        #line 224 "gpu_device_shm.impala"
        if (_30079) goto l30080; else goto l30111;
    l30111: ;
        #line 228 "gpu_device_shm.impala"
        goto l30108;
    l30080: ;
        #line 224 "gpu_device_shm.impala"
        bool _30082;
        _30082 = img_index_x_30078 < _30081;
        #line 224 "gpu_device_shm.impala"
        if (_30082) goto l30083; else goto l30110;
    l30110: ;
        #line 228 "gpu_device_shm.impala"
        goto l30108;
    l30083: ;
        #line 224 "gpu_device_shm.impala"
        if (_30087) goto l30088; else goto l30109;
    l30109: ;
        #line 228 "gpu_device_shm.impala"
        goto l30108;
    l30088: ;
        #line 224 "gpu_device_shm.impala"
        if (_30091) goto l30092; else goto l30107;
    l30107: ;
        #line 228 "gpu_device_shm.impala"
        goto l30108;
    l30108: ;
        #line 40 "gpu_device_shm.impala"
        goto l30093;
    l30092: ;
        #line 226 "gpu_device_shm.impala"
        int _30103;
        _30103 = _30102 + shm_index_x_30070;
        #line 226 "gpu_device_shm.impala"
        double* _30104;
        _30104 = reserve_shared_30050 + _30103;
        #line 227 "gpu_device_shm.impala"
        int _30097;
        _30097 = _30096 + img_index_x_30078;
        #line 227 "gpu_device_shm.impala"
        double* _30098;
        _30098 = _25297_30005 + _30097;
        #line 227 "gpu_device_shm.impala"
        double _30099;
        _30099 = *_30098;
        #line 227 "gpu_device_shm.impala"
        double _30105;
        _30105 = _30099;
        #line 226 "gpu_device_shm.impala"
        *_30104 = _30105;
        #line 40 "gpu_device_shm.impala"
        goto l30093;
    l30093: ;
        #line 42 "gpu_device_shm.impala"
        int _30095;
        _30095 = _30061 + bdimx_30012;
        #line 38 "gpu_device_shm.impala"
        p_30061 = _30095;
        goto l30060;
}

__global__ __launch_bounds__ (128 * 1 * 1) void lambda_25546(struct_image_6326 _25549_30238, struct_filter_6332 _25550_30239, double* _25551_30240, double* _25552_30241, struct_Buffer_6327 _25553_30242) {
    __shared__ double ds_img[134][7];
    int  bdimx_30245;
    int pbdimx_30245;
    int  bdimy_30248;
    int pbdimy_30248;
    int  bidx_30251;
    int pbidx_30251;
    int  bidy_30254;
    int pbidy_30254;
    int  tidx_30257;
    int ptidx_30257;
    int  tidy_30260;
    int ptidy_30260;
    double*  reserve_shared_30263;
    double* preserve_shared_30263;
    int  _30271;
    int p_30271;
    int  _30324;
    int p_30324;
    int  _30387;
    int p_30387;
    double  sum_30389;
    double psum_30389;
    int  _30329;
    int p_30329;
    #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
    bdimx_30245 = blockDim_x();
    pbdimx_30245 = bdimx_30245;
    l30243: ;
        bdimx_30245 = pbdimx_30245;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        bdimy_30248 = blockDim_y();
        pbdimy_30248 = bdimy_30248;
    l30246: ;
        bdimy_30248 = pbdimy_30248;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        bidx_30251 = blockIdx_x();
        pbidx_30251 = bidx_30251;
    l30249: ;
        bidx_30251 = pbidx_30251;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        bidy_30254 = blockIdx_y();
        pbidy_30254 = bidy_30254;
    l30252: ;
        bidy_30254 = pbidy_30254;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        tidx_30257 = threadIdx_x();
        ptidx_30257 = tidx_30257;
    l30255: ;
        tidx_30257 = ptidx_30257;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        tidy_30260 = threadIdx_y();
        ptidy_30260 = tidy_30260;
    l30258: ;
        tidy_30260 = ptidy_30260;
        #line 201 "gpu_device_shm.impala"
        __shared__ double reserver_reserve_shared_30263[938];
        preserve_shared_30263 = reserver_reserve_shared_30263;
    l30261: ;
        reserve_shared_30263 = preserve_shared_30263;
        #line 187 "gpu_device_shm.impala"
        int _30284;
        _30284 = bidx_30251 * bdimx_30245;
        #line 11 "main.impala"
        int _30298;
        _30298 = _25549_30238.e3;
        #line 187 "gpu_device_shm.impala"
        int gid_x_30285;
        gid_x_30285 = _30284 + tidx_30257;
        #line 6 "gaussian.impala"
        int _30264;
        _30264 = _25550_30239.e2;
        #line 191 "gpu_device_shm.impala"
        int _30273;
        _30273 = _25550_30239.e1;
        #line 189 "gpu_device_shm.impala"
        int _30293;
        _30293 = bidy_30254 * bdimy_30248;
        #line 11 "main.impala"
        int _30290;
        _30290 = _25549_30238.e2;
        #line 189 "gpu_device_shm.impala"
        int gid_y_30294;
        gid_y_30294 = _30293 + tidy_30260;
        #line 191 "gpu_device_shm.impala"
        int extend_width_30274;
        extend_width_30274 = _30273 / 2;
        #line 6 "gaussian.impala"
        int v_anchor_30265;
        v_anchor_30265 = _30264 / 2;
        #line 195 "gpu_device_shm.impala"
        int _30275;
        _30275 = 2 * extend_width_30274;
        #line 197 "gpu_device_shm.impala"
        int _30266;
        _30266 = 2 * v_anchor_30265;
        #line 195 "gpu_device_shm.impala"
        int shm_dimx_30276;
        shm_dimx_30276 = bdimx_30245 + _30275;
        #line 197 "gpu_device_shm.impala"
        int shm_dimy_30267;
        shm_dimy_30267 = bdimy_30248 + _30266;
        #line 38 "gpu_device_shm.impala"
        bool _30268;
        _30268 = 0 < shm_dimy_30267;
        #line 38 "gpu_device_shm.impala"
        if (_30268) goto l30269; else goto l30438;
    l30438: ;
        #line 231 "gpu_device_shm.impala"
        goto l30373;
    l30269: ;
        #line 226 "gpu_device_shm.impala"
        int _30309;
        _30309 = tidy_30260 * shm_dimx_30276;
        #line 219 "gpu_device_shm.impala"
        int _30286;
        _30286 = gid_x_30285 - extend_width_30274;
        #line 221 "gpu_device_shm.impala"
        int _30295;
        _30295 = gid_y_30294 - v_anchor_30265;
        #line 217 "gpu_device_shm.impala"
        bool _30282;
        _30282 = tidy_30260 < shm_dimy_30267;
        #line 224 "gpu_device_shm.impala"
        bool _30296;
        _30296 = 0 <= _30295;
        #line 224 "gpu_device_shm.impala"
        bool _30299;
        _30299 = _30295 < _30298;
        #line 227 "gpu_device_shm.impala"
        int _30304;
        _30304 = _30295 * _30290;
        #line 38 "gpu_device_shm.impala"
        p_30271 = 0;
        goto l30270;
    l30270: ;
        _30271 = p_30271;
        #line 38 "gpu_device_shm.impala"
        bool _30277;
        _30277 = _30271 < shm_dimx_30276;
        #line 38 "gpu_device_shm.impala"
        if (_30277) goto l30278; else goto l30322;
    l30322: ;
        #line 38 "gpu_device_shm.impala"
        p_30324 = bdimy_30248;
        goto l30323;
    l30323: ;
        _30324 = p_30324;
        #line 38 "gpu_device_shm.impala"
        bool _30326;
        _30326 = _30324 < shm_dimy_30267;
        #line 38 "gpu_device_shm.impala"
        if (_30326) goto l30327; else goto l30372;
    l30372: ;
        #line 231 "gpu_device_shm.impala"
        goto l30373;
    l30373: ;
        #line 1 "/home/rafael/repositories/anydsl/runtime/platforms/intrinsics_cuda.impala"
        __syncthreads();
    l30375: ;
        #line 237 "gpu_device_shm.impala"
        bool _30377;
        _30377 = gid_x_30285 < _30290;
        #line 237 "gpu_device_shm.impala"
        if (_30377) goto l30378; else goto l30437;
    l30437: ;
        #line 240 "gpu_device_shm.impala"
        goto l30436;
    l30378: ;
        #line 237 "gpu_device_shm.impala"
        bool _30379;
        _30379 = gid_y_30294 < _30298;
        #line 237 "gpu_device_shm.impala"
        if (_30379) goto l30380; else goto l30435;
    l30435: ;
        #line 240 "gpu_device_shm.impala"
        goto l30436;
    l30436: ;
        return ;
    l30380: ;
        #line 204 "gpu_device_shm.impala"
        int _30402;
        _30402 = v_anchor_30265 - _30293;
        #line 64 "gpu_device_shm.impala"
        char* _30416;
        _30416 = _25553_30242.e1;
        #line 203 "gpu_device_shm.impala"
        int _30406;
        _30406 = extend_width_30274 - _30284;
        #line 64 "gpu_device_shm.impala"
        int _30418;
        _30418 = gid_y_30294 * _30290;
        #line 64 "gpu_device_shm.impala"
        int _30419;
        _30419 = _30418 + gid_x_30285;
        #line 64 "gpu_device_shm.impala"
        double* _30417;
        union { double* dst; char* src; } u_30417;
        u_30417.src = _30416;
        _30417 = u_30417.dst;
        #line 39 "gaussian.impala"
        bool _30381;
        _30381 = v_anchor_30265 <= gid_y_30294;
        #line 64 "gpu_device_shm.impala"
        double* _30420;
        _30420 = _30417 + _30419;
        #line 39 "gaussian.impala"
        if (_30381) goto l30382; else goto l30434;
    l30434: ;
        #line 49 "gaussian.impala"
        goto l30424;
    l30382: ;
        #line 39 "gaussian.impala"
        int _30383;
        _30383 = _30298 - v_anchor_30265;
        #line 39 "gaussian.impala"
        bool _30384;
        _30384 = gid_y_30294 < _30383;
        #line 39 "gaussian.impala"
        if (_30384) goto l30385; else goto l30423;
    l30423: ;
        #line 49 "gaussian.impala"
        goto l30424;
    l30424: ;
        #line 72 "gpu_device_shm.impala"
        int _30425;
        _30425 = gid_y_30294 + _30402;
        #line 72 "gpu_device_shm.impala"
        int _30426;
        _30426 = _30425 * shm_dimx_30276;
        #line 72 "gpu_device_shm.impala"
        int _30427;
        _30427 = _30426 + gid_x_30285;
        #line 72 "gpu_device_shm.impala"
        int _30428;
        _30428 = _30427 + _30406;
        #line 72 "gpu_device_shm.impala"
        double* _30429;
        _30429 = reserve_shared_30263 + _30428;
        #line 72 "gpu_device_shm.impala"
        double _30430;
        _30430 = *_30429;
        #line 72 "gpu_device_shm.impala"
        double _30432;
        _30432 = _30430;
        #line 64 "gpu_device_shm.impala"
        *_30420 = _30432;
        return ;
    l30385: ;
        #line 41 "gaussian.impala"
        int _30422;
        _30422 = 0 - v_anchor_30265;
        #line 77 "gpu_device_shm.impala"
        struct_Buffer_6327 _30394;
        _30394 = _25550_30239.e0;
        #line 41 "gaussian.impala"
        int _30390;
        _30390 = 1 + v_anchor_30265;
        #line 77 "gpu_device_shm.impala"
        char* _30395;
        _30395 = _30394.e1;
        #line 77 "gpu_device_shm.impala"
        double* _30396;
        union { double* dst; char* src; } u_30396;
        u_30396.src = _30395;
        _30396 = u_30396.dst;
        #line 27 "gpu_device_shm.impala"
        p_30387 = _30422;
        psum_30389 = 0.000000e+00;
        goto l30386;
    l30386: ;
        _30387 = p_30387;
        sum_30389 = psum_30389;
        #line 27 "gpu_device_shm.impala"
        bool _30391;
        _30391 = _30387 < _30390;
        #line 27 "gpu_device_shm.impala"
        if (_30391) goto l30392; else goto l30415;
    l30415: ;
        #line 64 "gpu_device_shm.impala"
        *_30420 = sum_30389;
        return ;
    l30392: ;
        #line 43 "gaussian.impala"
        int _30401;
        _30401 = gid_y_30294 + _30387;
        #line 43 "gaussian.impala"
        int _30397;
        _30397 = _30387 + v_anchor_30265;
        #line 31 "gpu_device_shm.impala"
        int _30393;
        _30393 = 1 + _30387;
        #line 72 "gpu_device_shm.impala"
        int _30403;
        _30403 = _30401 + _30402;
        #line 76 "gpu_device_shm.impala"
        double* i_30398;
        i_30398 = _30396 + _30397;
        #line 72 "gpu_device_shm.impala"
        int _30404;
        _30404 = _30403 * shm_dimx_30276;
        #line 77 "gpu_device_shm.impala"
        double _30399;
        _30399 = *i_30398;
        #line 72 "gpu_device_shm.impala"
        int _30405;
        _30405 = _30404 + gid_x_30285;
        #line 77 "gpu_device_shm.impala"
        double _30411;
        _30411 = _30399;
        #line 72 "gpu_device_shm.impala"
        int _30407;
        _30407 = _30405 + _30406;
        #line 72 "gpu_device_shm.impala"
        double* _30408;
        _30408 = reserve_shared_30263 + _30407;
        #line 72 "gpu_device_shm.impala"
        double _30409;
        _30409 = *_30408;
        #line 72 "gpu_device_shm.impala"
        double _30412;
        _30412 = _30409;
        #line 43 "gaussian.impala"
        double _30413;
        _30413 = _30411 * _30412;
        #line 43 "gaussian.impala"
        double _30414;
        _30414 = sum_30389 + _30413;
        #line 27 "gpu_device_shm.impala"
        p_30387 = _30393;
        psum_30389 = _30414;
        goto l30386;
    l30327: ;
        #line 221 "gpu_device_shm.impala"
        int img_index_y_30344;
        img_index_y_30344 = _30295 + _30324;
        #line 224 "gpu_device_shm.impala"
        bool _30347;
        _30347 = img_index_y_30344 < _30298;
        #line 224 "gpu_device_shm.impala"
        bool _30345;
        _30345 = 0 <= img_index_y_30344;
        #line 214 "gpu_device_shm.impala"
        int shm_index_y_30336;
        shm_index_y_30336 = tidy_30260 + _30324;
        #line 227 "gpu_device_shm.impala"
        int _30352;
        _30352 = img_index_y_30344 * _30290;
        #line 217 "gpu_device_shm.impala"
        bool _30337;
        _30337 = shm_index_y_30336 < shm_dimy_30267;
        #line 226 "gpu_device_shm.impala"
        int _30357;
        _30357 = shm_index_y_30336 * shm_dimx_30276;
        #line 38 "gpu_device_shm.impala"
        p_30329 = 0;
        goto l30328;
    l30328: ;
        _30329 = p_30329;
        #line 38 "gpu_device_shm.impala"
        bool _30331;
        _30331 = _30329 < shm_dimx_30276;
        #line 38 "gpu_device_shm.impala"
        if (_30331) goto l30332; else goto l30370;
    l30370: ;
        #line 42 "gpu_device_shm.impala"
        int _30371;
        _30371 = _30324 + bdimy_30248;
        #line 38 "gpu_device_shm.impala"
        p_30324 = _30371;
        goto l30323;
    l30332: ;
        #line 212 "gpu_device_shm.impala"
        int shm_index_x_30333;
        shm_index_x_30333 = tidx_30257 + _30329;
        #line 217 "gpu_device_shm.impala"
        bool _30334;
        _30334 = shm_index_x_30333 < shm_dimx_30276;
        #line 217 "gpu_device_shm.impala"
        if (_30334) goto l30335; else goto l30369;
    l30369: ;
        #line 229 "gpu_device_shm.impala"
        goto l30368;
    l30335: ;
        #line 217 "gpu_device_shm.impala"
        if (_30337) goto l30338; else goto l30367;
    l30367: ;
        #line 229 "gpu_device_shm.impala"
        goto l30368;
    l30368: ;
        #line 40 "gpu_device_shm.impala"
        goto l30349;
    l30338: ;
        #line 219 "gpu_device_shm.impala"
        int img_index_x_30339;
        img_index_x_30339 = _30286 + _30329;
        #line 224 "gpu_device_shm.impala"
        bool _30340;
        _30340 = 0 <= img_index_x_30339;
        #line 224 "gpu_device_shm.impala"
        if (_30340) goto l30341; else goto l30366;
    l30366: ;
        #line 228 "gpu_device_shm.impala"
        goto l30363;
    l30341: ;
        #line 224 "gpu_device_shm.impala"
        bool _30342;
        _30342 = img_index_x_30339 < _30290;
        #line 224 "gpu_device_shm.impala"
        if (_30342) goto l30343; else goto l30365;
    l30365: ;
        #line 228 "gpu_device_shm.impala"
        goto l30363;
    l30343: ;
        #line 224 "gpu_device_shm.impala"
        if (_30345) goto l30346; else goto l30364;
    l30364: ;
        #line 228 "gpu_device_shm.impala"
        goto l30363;
    l30346: ;
        #line 224 "gpu_device_shm.impala"
        if (_30347) goto l30348; else goto l30362;
    l30362: ;
        #line 228 "gpu_device_shm.impala"
        goto l30363;
    l30363: ;
        #line 40 "gpu_device_shm.impala"
        goto l30349;
    l30348: ;
        #line 227 "gpu_device_shm.impala"
        int _30353;
        _30353 = _30352 + img_index_x_30339;
        #line 227 "gpu_device_shm.impala"
        double* _30354;
        _30354 = _25551_30240 + _30353;
        #line 227 "gpu_device_shm.impala"
        double _30355;
        _30355 = *_30354;
        #line 226 "gpu_device_shm.impala"
        int _30358;
        _30358 = _30357 + shm_index_x_30333;
        #line 227 "gpu_device_shm.impala"
        double _30360;
        _30360 = _30355;
        #line 226 "gpu_device_shm.impala"
        double* _30359;
        _30359 = reserve_shared_30263 + _30358;
        #line 226 "gpu_device_shm.impala"
        *_30359 = _30360;
        #line 40 "gpu_device_shm.impala"
        goto l30349;
    l30349: ;
        #line 42 "gpu_device_shm.impala"
        int _30351;
        _30351 = _30329 + bdimx_30245;
        #line 38 "gpu_device_shm.impala"
        p_30329 = _30351;
        goto l30328;
    l30278: ;
        #line 212 "gpu_device_shm.impala"
        int shm_index_x_30279;
        shm_index_x_30279 = tidx_30257 + _30271;
        #line 217 "gpu_device_shm.impala"
        bool _30280;
        _30280 = shm_index_x_30279 < shm_dimx_30276;
        #line 217 "gpu_device_shm.impala"
        if (_30280) goto l30281; else goto l30321;
    l30321: ;
        #line 229 "gpu_device_shm.impala"
        goto l30320;
    l30281: ;
        #line 217 "gpu_device_shm.impala"
        if (_30282) goto l30283; else goto l30319;
    l30319: ;
        #line 229 "gpu_device_shm.impala"
        goto l30320;
    l30320: ;
        #line 40 "gpu_device_shm.impala"
        goto l30301;
    l30283: ;
        #line 219 "gpu_device_shm.impala"
        int img_index_x_30287;
        img_index_x_30287 = _30286 + _30271;
        #line 224 "gpu_device_shm.impala"
        bool _30288;
        _30288 = 0 <= img_index_x_30287;
        #line 224 "gpu_device_shm.impala"
        if (_30288) goto l30289; else goto l30318;
    l30318: ;
        #line 228 "gpu_device_shm.impala"
        goto l30315;
    l30289: ;
        #line 224 "gpu_device_shm.impala"
        bool _30291;
        _30291 = img_index_x_30287 < _30290;
        #line 224 "gpu_device_shm.impala"
        if (_30291) goto l30292; else goto l30317;
    l30317: ;
        #line 228 "gpu_device_shm.impala"
        goto l30315;
    l30292: ;
        #line 224 "gpu_device_shm.impala"
        if (_30296) goto l30297; else goto l30316;
    l30316: ;
        #line 228 "gpu_device_shm.impala"
        goto l30315;
    l30297: ;
        #line 224 "gpu_device_shm.impala"
        if (_30299) goto l30300; else goto l30314;
    l30314: ;
        #line 228 "gpu_device_shm.impala"
        goto l30315;
    l30315: ;
        #line 40 "gpu_device_shm.impala"
        goto l30301;
    l30300: ;
        #line 226 "gpu_device_shm.impala"
        int _30310;
        _30310 = _30309 + shm_index_x_30279;
        #line 227 "gpu_device_shm.impala"
        int _30305;
        _30305 = _30304 + img_index_x_30287;
        #line 226 "gpu_device_shm.impala"
        double* _30311;
        _30311 = reserve_shared_30263 + _30310;
        #line 227 "gpu_device_shm.impala"
        double* _30306;
        _30306 = _25551_30240 + _30305;
        #line 227 "gpu_device_shm.impala"
        double _30307;
        _30307 = *_30306;
        #line 227 "gpu_device_shm.impala"
        double _30312;
        _30312 = _30307;
        #line 226 "gpu_device_shm.impala"
        *_30311 = _30312;
        #line 40 "gpu_device_shm.impala"
        goto l30301;
    l30301: ;
        #line 42 "gpu_device_shm.impala"
        int _30303;
        _30303 = _30271 + bdimx_30245;
        #line 38 "gpu_device_shm.impala"
        p_30271 = _30303;
        goto l30270;
}

}
	.text
	.file	"main_shm.hip"
	.globl	__device_stub__lambda_25292     # -- Begin function __device_stub__lambda_25292
	.type	__device_stub__lambda_25292,@function
__device_stub__lambda_25292:            # @__device_stub__lambda_25292
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	64(%rsp), %rax
	movl	%edi, (%rax)
	movq	%rsi, 8(%rax)
	leaq	24(%rsp), %rcx
	movq	%rdx, (%rcx)
	leaq	160(%rsp), %rdx
	leaq	80(%rsp), %rbx
	movq	%rdx, (%rbx)
	movq	%rax, 8(%rbx)
	movq	%rcx, 16(%rbx)
	leaq	200(%rsp), %rax
	movq	%rax, 24(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$lambda_25292, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	__device_stub__lambda_25292, .Lfunc_end0-__device_stub__lambda_25292
	.cfi_endproc
                                        # -- End function
	.globl	__device_stub__lambda_25546     # -- Begin function __device_stub__lambda_25546
	.type	__device_stub__lambda_25546,@function
__device_stub__lambda_25546:            # @__device_stub__lambda_25546
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$128, %rsp
	.cfi_def_cfa_offset 176
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	64(%rsp), %rax
	movl	%edx, (%rax)
	movq	%rcx, 8(%rax)
	leaq	24(%rsp), %rcx
	movq	%rdi, (%rcx)
	leaq	16(%rsp), %rdx
	movq	%rsi, (%rdx)
	leaq	176(%rsp), %rsi
	leaq	80(%rsp), %rbx
	movq	%rsi, (%rbx)
	leaq	216(%rsp), %rsi
	movq	%rsi, 8(%rbx)
	movq	%rcx, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rax, 32(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%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	$lambda_25546, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$144, %rsp
	.cfi_adjust_cfa_offset -144
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	__device_stub__lambda_25546, .Lfunc_end1-__device_stub__lambda_25546
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB2_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB2_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$lambda_25292, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$lambda_25546, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end2:
	.size	__hip_module_ctor, .Lfunc_end2-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB3_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB3_2:
	retq
.Lfunc_end3:
	.size	__hip_module_dtor, .Lfunc_end3-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	lambda_25292,@object            # @lambda_25292
	.section	.rodata,"a",@progbits
	.globl	lambda_25292
	.p2align	3, 0x0
lambda_25292:
	.quad	__device_stub__lambda_25292
	.size	lambda_25292, 8

	.type	lambda_25546,@object            # @lambda_25546
	.globl	lambda_25546
	.p2align	3, 0x0
lambda_25546:
	.quad	__device_stub__lambda_25546
	.size	lambda_25546, 8

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

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"lambda_25546"
	.size	.L__unnamed_2, 13

	.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 __device_stub__lambda_25292
	.addrsig_sym __device_stub__lambda_25546
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym lambda_25292
	.addrsig_sym lambda_25546
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	lambda_25292            ; -- Begin function lambda_25292
	.globl	lambda_25292
	.p2align	8
	.type	lambda_25292,@function
lambda_25292:                           ; @lambda_25292
; %bb.0:
	s_clause 0x3
	s_load_b32 s2, s[0:1], 0x64
	s_load_b64 s[12:13], s[0:1], 0x50
	s_load_b64 s[6:7], s[0:1], 0x20
	s_load_b128 s[8:11], s[0:1], 0x30
	v_bfe_u32 v2, v0, 10, 10
	v_and_b32_e32 v6, 0x3ff, v0
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s18, s2, 16
	s_and_b32 s19, s2, 0xffff
	s_lshr_b32 s2, s13, 31
	s_lshr_b32 s3, s12, 31
	s_add_i32 s2, s13, s2
	v_mad_u64_u32 v[0:1], null, s15, s18, v[2:3]
	s_mul_i32 s14, s14, s19
	s_add_i32 s3, s12, s3
	s_and_b32 s20, s2, -2
	v_add_nc_u32_e32 v1, s14, v6
	s_and_b32 s21, s3, -2
	s_add_i32 s22, s20, s18
	s_ashr_i32 s16, s2, 1
	s_ashr_i32 s5, s3, 1
	s_add_i32 s17, s21, s19
	s_cmp_lt_i32 s22, 1
	s_cbranch_scc1 .LBB0_16
; %bb.1:
	v_subrev_nc_u32_e32 v7, s16, v0
	v_subrev_nc_u32_e32 v3, s5, v1
	s_cmp_gt_i32 s17, 0
	s_mov_b32 s24, 0
	s_cselect_b32 s23, -1, 0
	s_cmp_lt_i32 s17, 1
	s_cbranch_scc1 .LBB0_7
; %bb.2:                                ; %.lr.ph
	v_mul_lo_u32 v8, v2, s17
	v_lshlrev_b32_e32 v9, 3, v6
	v_mad_u64_u32 v[4:5], null, v7, s6, v[3:4]
	v_cmp_le_u32_e64 s3, s22, v2
	v_cmp_gt_i32_e32 vcc_lo, 0, v7
	v_cmp_le_i32_e64 s2, s7, v7
	s_lshl_b32 s25, s19, 3
	v_lshl_add_u32 v5, v8, 3, v9
	s_xor_b32 s26, s3, -1
.LBB0_3:                                ; =>This Inner Loop Header: Depth=1
	v_add_nc_u32_e32 v8, s24, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_gt_i32_e64 s3, s17, v8
	s_and_b32 s3, s26, s3
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s27, s3
	s_cbranch_execz .LBB0_6
; %bb.4:                                ;   in Loop: Header=BB0_3 Depth=1
	v_add_nc_u32_e32 v8, s24, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_gt_i32_e64 s3, 0, v8
	v_cmp_le_i32_e64 s4, s6, v8
	s_or_b32 s3, s3, s4
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s3, s3, vcc_lo
	s_or_b32 s3, s3, s2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_xor_b32 s3, s3, -1
	s_and_b32 exec_lo, exec_lo, s3
	s_cbranch_execz .LBB0_6
; %bb.5:                                ;   in Loop: Header=BB0_3 Depth=1
	v_add_nc_u32_e32 v8, s24, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v9, 31, v8
	v_lshlrev_b64 v[8:9], 3, v[8:9]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v8, s3, s10, v8
	v_add_co_ci_u32_e64 v9, s3, s11, v9, s3
	global_load_b64 v[8:9], v[8:9], off
	s_waitcnt vmcnt(0)
	ds_store_b64 v5, v[8:9]
.LBB0_6:                                ;   in Loop: Header=BB0_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s27
	v_add_nc_u32_e32 v5, s25, v5
	s_add_i32 s3, s24, s19
	s_cmp_ge_i32 s24, s21
	s_mov_b32 s24, s3
	s_cbranch_scc0 .LBB0_3
.LBB0_7:                                ; %.preheader
	s_cmp_lt_i32 s13, 2
	s_cbranch_scc1 .LBB0_16
; %bb.8:                                ; %.lr.ph256.preheader
	s_add_i32 s15, s15, 1
	v_lshlrev_b32_e32 v10, 3, v6
	v_mad_u64_u32 v[4:5], null, s15, s18, v[2:3]
	v_add_nc_u32_e32 v5, s18, v2
	s_mul_i32 s2, s17, s18
	s_mul_i32 s13, s6, s18
	s_lshl_b32 s15, s2, 3
	s_lshl_b32 s24, s19, 3
	v_mul_lo_u32 v8, v5, s17
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_subrev_nc_u32_e32 v9, s16, v4
	s_mov_b32 s25, 0
	s_mov_b32 s26, s18
	v_mad_u64_u32 v[4:5], null, s6, v9, v[3:4]
	s_delay_alu instid0(VALU_DEP_3)
	v_lshl_add_u32 v5, v8, 3, v10
.LBB0_9:                                ; %.lr.ph256
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_11 Depth 2
	s_and_not1_b32 vcc_lo, exec_lo, s23
	s_cbranch_vccnz .LBB0_15
; %bb.10:                               ; %.lr.ph254
                                        ;   in Loop: Header=BB0_9 Depth=1
	v_add_nc_u32_e32 v9, s26, v2
	v_add_nc_u32_e32 v8, s26, v7
	s_mov_b32 s28, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_cmp_le_i32_e64 s3, s22, v9
	v_mov_b32_e32 v9, v5
	v_cmp_gt_i32_e32 vcc_lo, 0, v8
	v_cmp_le_i32_e64 s2, s7, v8
	v_add_nc_u32_e32 v8, s25, v4
	s_xor_b32 s27, s3, -1
.LBB0_11:                               ;   Parent Loop BB0_9 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	v_add_nc_u32_e32 v10, s28, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_gt_i32_e64 s3, s17, v10
	s_and_b32 s3, s3, s27
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s29, s3
	s_cbranch_execz .LBB0_14
; %bb.12:                               ;   in Loop: Header=BB0_11 Depth=2
	v_add_nc_u32_e32 v10, s28, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_gt_i32_e64 s3, 0, v10
	v_cmp_le_i32_e64 s4, s6, v10
	s_or_b32 s3, s3, s4
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s3, s3, vcc_lo
	s_or_b32 s3, s3, s2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_xor_b32 s3, s3, -1
	s_and_b32 exec_lo, exec_lo, s3
	s_cbranch_execz .LBB0_14
; %bb.13:                               ;   in Loop: Header=BB0_11 Depth=2
	v_add_nc_u32_e32 v10, s28, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v11, 31, v10
	v_lshlrev_b64 v[10:11], 3, v[10:11]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v10, s3, s10, v10
	v_add_co_ci_u32_e64 v11, s3, s11, v11, s3
	global_load_b64 v[10:11], v[10:11], off
	s_waitcnt vmcnt(0)
	ds_store_b64 v9, v[10:11]
.LBB0_14:                               ;   in Loop: Header=BB0_11 Depth=2
	s_or_b32 exec_lo, exec_lo, s29
	v_add_nc_u32_e32 v9, s24, v9
	s_add_i32 s3, s28, s19
	s_cmp_ge_i32 s28, s21
	s_mov_b32 s28, s3
	s_cbranch_scc0 .LBB0_11
.LBB0_15:                               ; %._crit_edge
                                        ;   in Loop: Header=BB0_9 Depth=1
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v5, s15, v5
	s_add_i32 s2, s26, s18
	s_add_i32 s25, s25, s13
	s_cmp_ge_i32 s26, s20
	s_mov_b32 s26, s2
	s_cbranch_scc0 .LBB0_9
.LBB0_16:                               ; %.loopexit
	v_cmp_gt_i32_e32 vcc_lo, s6, v1
	v_cmp_gt_i32_e64 s2, s7, v0
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_and_b32 s2, vcc_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s3, s2
	s_cbranch_execz .LBB0_24
; %bb.17:
	v_add_nc_u32_e32 v2, s16, v2
	s_sub_i32 s2, s6, s5
	v_cmp_gt_i32_e32 vcc_lo, s5, v1
	v_cmp_le_i32_e64 s2, s2, v1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v4, s17, v2
                                        ; implicit-def: $vgpr2_vgpr3
	s_or_b32 s2, vcc_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_saveexec_b32 s3, s2
	s_xor_b32 s2, exec_lo, s3
	s_cbranch_execz .LBB0_19
; %bb.18:
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	v_lshlrev_b32_e32 v2, 3, v4
	v_lshlrev_b32_e32 v3, 3, v1
	s_sub_i32 s3, s5, s14
                                        ; implicit-def: $vgpr6
                                        ; implicit-def: $vgpr4
	s_lshl_b32 s3, s3, 3
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_add3_u32 v2, v2, v3, s3
	ds_load_b64 v[2:3], v2
.LBB0_19:                               ; %Flow287
	s_and_not1_saveexec_b32 s4, s2
	s_cbranch_execz .LBB0_23
; %bb.20:
	s_waitcnt lgkmcnt(0)
	v_mov_b32_e32 v2, 0
	v_mov_b32_e32 v3, 0
	s_cmp_lt_i32 s12, -1
	s_cbranch_scc1 .LBB0_23
; %bb.21:                               ; %.lr.ph261
	s_load_b64 s[0:1], s[0:1], 0x48
	v_dual_mov_b32 v2, 0 :: v_dual_lshlrev_b32 v5, 3, v6
	s_abs_i32 s2, s5
	v_mov_b32_e32 v3, 0
	s_add_i32 s5, s5, s2
	s_delay_alu instid0(VALU_DEP_2)
	v_lshl_add_u32 v4, v4, 3, v5
	s_add_i32 s5, s5, 1
	s_mov_b32 s2, 0
.LBB0_22:                               ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_ashr_i32 s3, s2, 31
	ds_load_b64 v[5:6], v4
	s_lshl_b64 s[10:11], s[2:3], 3
	v_add_nc_u32_e32 v4, 8, v4
	s_waitcnt lgkmcnt(0)
	s_add_u32 s10, s0, s10
	s_addc_u32 s11, s1, s11
	s_add_i32 s2, s2, 1
	s_load_b64 s[10:11], s[10:11], 0x0
	s_cmp_eq_u32 s5, s2
	s_waitcnt lgkmcnt(0)
	v_fma_f64 v[2:3], s[10:11], v[5:6], v[2:3]
	s_cbranch_scc0 .LBB0_22
.LBB0_23:                               ; %.sink.split
	s_or_b32 exec_lo, exec_lo, s4
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[4:5], null, v0, s6, v[1:2]
	v_ashrrev_i32_e32 v5, 31, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 3, v[4:5]
	v_add_co_u32 v0, vcc_lo, s8, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s9, v1, vcc_lo
	global_store_b64 v[0:1], v[2:3], off
.LBB0_24:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel lambda_25292
		.amdhsa_group_segment_fixed_size 7504
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 344
		.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 12
		.amdhsa_next_free_sgpr 30
		.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	lambda_25292, .Lfunc_end0-lambda_25292
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1084
; NumSgprs: 32
; NumVgprs: 12
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 7504 bytes/workgroup (compile time only)
; SGPRBlocks: 3
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 32
; NumVGPRsForWavesPerEU: 12
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.protected	lambda_25546            ; -- Begin function lambda_25546
	.globl	lambda_25546
	.p2align	8
	.type	lambda_25546,@function
lambda_25546:                           ; @lambda_25546
; %bb.0:
	s_clause 0x2
	s_load_b32 s2, s[0:1], 0x6c
	s_load_b64 s[4:5], s[0:1], 0x38
	s_load_b64 s[6:7], s[0:1], 0x20
	v_bfe_u32 v2, v0, 10, 10
	v_and_b32_e32 v6, 0x3ff, v0
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s16, s2, 16
	s_lshr_b32 s3, s5, 31
	s_and_b32 s10, s2, 0xffff
	s_lshr_b32 s2, s4, 31
	v_mad_u64_u32 v[0:1], null, s15, s16, v[2:3]
	s_add_i32 s3, s5, s3
	s_mul_i32 s14, s14, s10
	s_add_i32 s2, s4, s2
	s_and_b32 s17, s3, -2
	v_add_nc_u32_e32 v1, s14, v6
	s_and_b32 s18, s2, -2
	s_add_i32 s19, s17, s16
	s_ashr_i32 s11, s2, 1
	s_ashr_i32 s12, s3, 1
	s_add_i32 s13, s18, s10
	s_cmp_lt_i32 s19, 1
	s_cbranch_scc1 .LBB1_16
; %bb.1:
	s_load_b64 s[8:9], s[0:1], 0x40
	v_subrev_nc_u32_e32 v3, s11, v1
	v_subrev_nc_u32_e32 v7, s12, v0
	v_lshlrev_b32_e32 v8, 3, v6
	s_cmp_gt_i32 s13, 0
	s_mov_b32 s21, 0
	s_cselect_b32 s20, -1, 0
	s_cmp_lt_i32 s13, 1
	s_cbranch_scc1 .LBB1_7
; %bb.2:                                ; %.lr.ph
	v_mul_lo_u32 v9, v2, s13
	v_mad_u64_u32 v[4:5], null, v7, s6, v[3:4]
	v_cmp_le_u32_e64 s3, s19, v2
	v_cmp_gt_i32_e32 vcc_lo, 0, v7
	v_cmp_le_i32_e64 s2, s7, v7
	s_lshl_b32 s22, s10, 3
	s_delay_alu instid0(VALU_DEP_3)
	s_xor_b32 s23, s3, -1
	v_lshl_add_u32 v5, v9, 3, v8
.LBB1_3:                                ; =>This Inner Loop Header: Depth=1
	v_add_nc_u32_e32 v9, s21, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_gt_i32_e64 s3, s13, v9
	s_and_b32 s3, s23, s3
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s24, s3
	s_cbranch_execz .LBB1_6
; %bb.4:                                ;   in Loop: Header=BB1_3 Depth=1
	v_add_nc_u32_e32 v9, s21, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_gt_i32_e64 s3, 0, v9
	v_cmp_le_i32_e64 s4, s6, v9
	s_or_b32 s3, s3, s4
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s3, s3, vcc_lo
	s_or_b32 s3, s3, s2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_xor_b32 s3, s3, -1
	s_and_b32 exec_lo, exec_lo, s3
	s_cbranch_execz .LBB1_6
; %bb.5:                                ;   in Loop: Header=BB1_3 Depth=1
	v_add_nc_u32_e32 v9, s21, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v10, 31, v9
	v_lshlrev_b64 v[9:10], 3, v[9:10]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v9, s3, s8, v9
	v_add_co_ci_u32_e64 v10, s3, s9, v10, s3
	global_load_b64 v[9:10], v[9:10], off
	s_waitcnt vmcnt(0)
	ds_store_b64 v5, v[9:10]
.LBB1_6:                                ;   in Loop: Header=BB1_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s24
	v_add_nc_u32_e32 v5, s22, v5
	s_add_i32 s3, s21, s10
	s_cmp_ge_i32 s21, s18
	s_mov_b32 s21, s3
	s_cbranch_scc0 .LBB1_3
.LBB1_7:                                ; %.preheader
	s_cmp_lt_i32 s5, 2
	s_cbranch_scc1 .LBB1_16
; %bb.8:                                ; %.lr.ph261.preheader
	s_add_i32 s15, s15, 1
	s_mul_i32 s2, s13, s16
	v_mad_u64_u32 v[4:5], null, s15, s16, v[2:3]
	v_add_nc_u32_e32 v5, s16, v2
	s_mul_i32 s15, s6, s16
	s_lshl_b32 s21, s2, 3
	s_lshl_b32 s22, s10, 3
	s_mov_b32 s23, 0
	v_mul_lo_u32 v9, v5, s13
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_subrev_nc_u32_e32 v10, s12, v4
	s_mov_b32 s24, s16
	v_mad_u64_u32 v[4:5], null, s6, v10, v[3:4]
	s_delay_alu instid0(VALU_DEP_3)
	v_lshl_add_u32 v5, v9, 3, v8
.LBB1_9:                                ; %.lr.ph261
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB1_11 Depth 2
	s_and_not1_b32 vcc_lo, exec_lo, s20
	s_cbranch_vccnz .LBB1_15
; %bb.10:                               ; %.lr.ph259
                                        ;   in Loop: Header=BB1_9 Depth=1
	v_add_nc_u32_e32 v9, s24, v2
	v_add_nc_u32_e32 v8, s24, v7
	s_mov_b32 s26, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_cmp_le_i32_e64 s3, s19, v9
	v_mov_b32_e32 v9, v5
	v_cmp_le_i32_e32 vcc_lo, s7, v8
	v_cmp_gt_i32_e64 s2, 0, v8
	v_add_nc_u32_e32 v8, s23, v4
	s_xor_b32 s25, s3, -1
.LBB1_11:                               ;   Parent Loop BB1_9 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	v_add_nc_u32_e32 v10, s26, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_gt_i32_e64 s3, s13, v10
	s_and_b32 s3, s3, s25
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s27, s3
	s_cbranch_execz .LBB1_14
; %bb.12:                               ;   in Loop: Header=BB1_11 Depth=2
	v_add_nc_u32_e32 v10, s26, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_gt_i32_e64 s3, 0, v10
	v_cmp_le_i32_e64 s4, s6, v10
	s_or_b32 s3, s3, s4
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s3, s3, s2
	s_or_b32 s3, s3, vcc_lo
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_xor_b32 s3, s3, -1
	s_and_b32 exec_lo, exec_lo, s3
	s_cbranch_execz .LBB1_14
; %bb.13:                               ;   in Loop: Header=BB1_11 Depth=2
	v_add_nc_u32_e32 v10, s26, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v11, 31, v10
	v_lshlrev_b64 v[10:11], 3, v[10:11]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v10, s3, s8, v10
	v_add_co_ci_u32_e64 v11, s3, s9, v11, s3
	global_load_b64 v[10:11], v[10:11], off
	s_waitcnt vmcnt(0)
	ds_store_b64 v9, v[10:11]
.LBB1_14:                               ;   in Loop: Header=BB1_11 Depth=2
	s_or_b32 exec_lo, exec_lo, s27
	v_add_nc_u32_e32 v9, s22, v9
	s_add_i32 s3, s26, s10
	s_cmp_ge_i32 s26, s18
	s_mov_b32 s26, s3
	s_cbranch_scc0 .LBB1_11
.LBB1_15:                               ; %._crit_edge
                                        ;   in Loop: Header=BB1_9 Depth=1
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v5, s21, v5
	s_add_i32 s2, s24, s16
	s_add_i32 s23, s23, s15
	s_cmp_ge_i32 s24, s17
	s_mov_b32 s24, s2
	s_cbranch_scc0 .LBB1_9
.LBB1_16:                               ; %.loopexit
	v_cmp_gt_i32_e32 vcc_lo, s6, v1
	v_cmp_gt_i32_e64 s2, s7, v0
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_and_b32 s2, vcc_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s3, s2
	s_cbranch_execz .LBB1_24
; %bb.17:
	s_load_b64 s[8:9], s[0:1], 0x58
	s_sub_i32 s2, s7, s12
	v_cmp_gt_i32_e32 vcc_lo, s12, v0
	v_cmp_le_i32_e64 s2, s2, v0
                                        ; implicit-def: $vgpr3_vgpr4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s2, vcc_lo, s2
	s_and_saveexec_b32 s3, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s2, exec_lo, s3
	s_cbranch_execz .LBB1_19
; %bb.18:
	v_add_nc_u32_e32 v2, s12, v2
	v_lshlrev_b32_e32 v3, 3, v1
	s_sub_i32 s3, s11, s14
                                        ; implicit-def: $vgpr6
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	s_lshl_b32 s3, s3, 3
	v_mul_lo_u32 v2, s13, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b32_e32 v2, 3, v2
	v_add3_u32 v2, v2, v3, s3
	ds_load_b64 v[3:4], v2
                                        ; implicit-def: $vgpr2
.LBB1_19:                               ; %Flow292
	s_and_not1_saveexec_b32 s4, s2
	s_cbranch_execz .LBB1_23
; %bb.20:
	s_waitcnt lgkmcnt(0)
	v_mov_b32_e32 v3, 0
	v_mov_b32_e32 v4, 0
	s_cmp_lt_i32 s5, -1
	s_cbranch_scc1 .LBB1_23
; %bb.21:                               ; %.lr.ph268
	s_load_b64 s[0:1], s[0:1], 0x30
	v_mul_lo_u32 v2, v2, s13
	v_add_lshl_u32 v3, v6, s11, 3
	s_abs_i32 s2, s12
	s_lshl_b32 s7, s11, 4
	s_add_i32 s5, s12, s2
	s_lshl_b32 s2, s10, 3
	s_add_i32 s5, s5, 1
	s_add_i32 s7, s7, s2
	v_lshl_add_u32 v2, v2, 3, v3
	v_mov_b32_e32 v3, 0
	v_mov_b32_e32 v4, 0
	s_mov_b32 s2, 0
.LBB1_22:                               ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_ashr_i32 s3, s2, 31
	ds_load_b64 v[5:6], v2
	s_lshl_b64 s[10:11], s[2:3], 3
	v_add_nc_u32_e32 v2, s7, v2
	s_waitcnt lgkmcnt(0)
	s_add_u32 s10, s0, s10
	s_addc_u32 s11, s1, s11
	s_add_i32 s2, s2, 1
	s_load_b64 s[10:11], s[10:11], 0x0
	s_cmp_eq_u32 s5, s2
	s_waitcnt lgkmcnt(0)
	v_fma_f64 v[3:4], s[10:11], v[5:6], v[3:4]
	s_cbranch_scc0 .LBB1_22
.LBB1_23:                               ; %.sink.split
	s_or_b32 exec_lo, exec_lo, s4
	v_mad_u64_u32 v[5:6], null, v0, s6, v[1:2]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v6, 31, v5
	v_lshlrev_b64 v[0:1], 3, v[5:6]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s8, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s9, v1, vcc_lo
	global_store_b64 v[0:1], v[3:4], off
.LBB1_24:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel lambda_25546
		.amdhsa_group_segment_fixed_size 7504
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 352
		.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 12
		.amdhsa_next_free_sgpr 28
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end1:
	.size	lambda_25546, .Lfunc_end1-lambda_25546
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1116
; NumSgprs: 30
; NumVgprs: 12
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 7504 bytes/workgroup (compile time only)
; SGPRBlocks: 3
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 30
; NumVGPRsForWavesPerEU: 12
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .offset:         0
        .size:           40
        .value_kind:     by_value
      - .offset:         40
        .size:           16
        .value_kind:     by_value
      - .address_space:  global
        .offset:         56
        .size:           8
        .value_kind:     global_buffer
      - .offset:         64
        .size:           24
        .value_kind:     by_value
      - .offset:         88
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         92
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         96
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         100
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         102
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         104
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         106
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         108
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         110
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         128
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         136
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         144
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         152
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 7504
    .kernarg_segment_align: 8
    .kernarg_segment_size: 344
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 128
    .name:           lambda_25292
    .private_segment_fixed_size: 0
    .sgpr_count:     32
    .sgpr_spill_count: 0
    .symbol:         lambda_25292.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     12
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .offset:         0
        .size:           40
        .value_kind:     by_value
      - .offset:         40
        .size:           24
        .value_kind:     by_value
      - .address_space:  global
        .offset:         64
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         72
        .size:           8
        .value_kind:     global_buffer
      - .offset:         80
        .size:           16
        .value_kind:     by_value
      - .offset:         96
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         100
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         104
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         108
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         110
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         112
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         114
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         116
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         118
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         136
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         144
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         152
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         160
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 7504
    .kernarg_segment_align: 8
    .kernarg_segment_size: 352
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 128
    .name:           lambda_25546
    .private_segment_fixed_size: 0
    .sgpr_count:     30
    .sgpr_spill_count: 0
    .symbol:         lambda_25546.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     12
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
