     const long int min_out_x = offset_x + last_x*shift_x;
     const long int min_out_y = offset_y + last_y*shift_y;
 
-    for(long int out_x = 0; out_x < length_x; out_x++)
-    {
-        if((out_x + min_out_x < 0) || (out_x + min_out_x >= output_size_x))
-        {
-            continue;
-        }
-
-        for(long int out_y = 0; out_y < length_y; out_y++)
-        {
-            if((out_y + min_out_y < 0) || (out_y + min_out_y >= output_size_y))
-            {
-                continue;
-            }
-
+    const long int min_x = (min_out_x < 0) ? -min_out_x : 0;
+    const long int min_y = (min_out_y < 0) ? -min_out_y : 0;
+
+    const long int max_x = (length_x + min_out_x > output_size_x) ? output_size_x - min_out_x : length_x;
+    const long int max_y = (length_y + min_out_y > output_size_y) ? output_size_y - min_out_y : length_y;
+
+    for(long int out_x = min_x; out_x < max_x; out_x++)
+    {
+        for(long int out_y = min_y; out_y < max_y; out_y++)
+        {
             for(size_t out_z = 0; out_z < output_size_z; out_z++)
             {
                 size_t i = (out_x + min_out_x)*output_size_y*output_size_z + (out_y + min_out_y)*output_size_z + out_z;

    const size_t last_y = blockIdx.y;
    const size_t last_z = blockIdx.z;

    const size_t n = last_x*gridDim.y*gridDim.z + last_y*gridDim.z + last_z;

    const long int min_out_x = offset_x + last_x*shift_x;
    const long int min_out_y = offset_y + last_y*shift_y;

    size_t inp_z = threadIdx.x;

    float_type tmp1, tmp2;
    float_type maxRes, maxMul;

    while(inp_z < input_size_z)
    {
        for(long int out_x = 0; out_x < length_x; out_x++)
        {
            if((out_x + min_out_x < 0) || (out_x + min_out_x >= output_size_x))
            {
                continue;
            }

            for(long int out_y = 0; out_y < length_y; out_y++)
            {
                if((out_y + min_out_y < 0) || (out_y + min_out_y >= output_size_y))
                {
                    continue;
                }

                for(size_t x_shift = 0; x_shift < filter_size_x; x_shift++)
                {
                    for(size_t y_shift = 0; y_shift < filter_size_y; y_shift++)

    const size_t last_y = blockIdx.y;
    const size_t last_z = blockIdx.z;

    const size_t inp_z = threadIdx.x;
    const size_t y_shift = threadIdx.y;
    const size_t x_shift = threadIdx.z;

    __shared__ size_t n;

    __shared__ long int min_out_x;
    __shared__ long int min_out_y;

    __shared__ long int min_x;
    __shared__ long int min_y;

    __shared__ long int max_x;
    __shared__ long int max_y;

    if(x_shift*filter_size_y*input_size_z + y_shift*input_size_z + inp_z == 0)
    {
        n = last_x*gridDim.y*gridDim.z + last_y*gridDim.z + last_z;

        min_out_x = offset_x + last_x*shift_x;
        min_out_y = offset_y + last_y*shift_y;

        min_x = (min_out_x < 0) ? -min_out_x : 0;
        min_y = (min_out_y < 0) ? -min_out_y : 0;

        max_x = (length_x + min_out_x > output_size_x) ? output_size_x - min_out_x : length_x;
        max_y = (length_y + min_out_y > output_size_y) ? output_size_y - min_out_y : length_y;
    }

    __syncthreads();

    float_type tmp1, tmp2;
    float_type maxRes, maxMul;

    if(x_shift*filter_size_y*input_size_z + y_shift*input_size_z + inp_z < filter_size_x*filter_size_y*input_size_z)
    {
        for(long int out_x = min_x; out_x < max_x; out_x++)
        {
            for(long int out_y = min_y; out_y < max_y; out_y++)
            {
                const long int x_val = out_x*stride_x + min_out_x*stride_x + x_shift - pad_x;
                const long int y_val = out_y*stride_y + min_out_y*stride_y + y_shift - pad_y;

