o
    :ήc                     @   sp   d dl Zd dlZdZe d  Zeejj	d  Z
ee
gZeeZeee Z					d
dd	ZdS )    Na	  
__kernel void resize_nearest(
    __global float *dst,
    __global const float *src,
    int w_src, int h_src,
    int w_dst, int h_dst,
    int depth
){
    int x_dst = get_global_id(0);
    int y_dst = get_global_id(1);

    if (x_dst >= w_dst || y_dst >= h_dst) return;

    int x_src = min(x_dst * w_src / w_dst, w_src - 1);
    int y_src = min(y_dst * h_src / h_dst, h_src - 1);

    __global       float *ptr_dst = dst + (x_dst + y_dst * w_dst) * depth;
    __global const float *ptr_src = src + (x_src + y_src * w_src) * depth;

    for (int channel = 0; channel < depth; channel++){
        ptr_dst[channel] = ptr_src[channel];
    }
}

__kernel void ml_iteration(
    __global       float *F,
    __global       float *B,
    __global const float *F_prev,
    __global const float *B_prev,
    __global const float *image,
    __global const float *alpha,
    int w,
    int h,
    float regularization
){
    int x = get_global_id(0);
    int y = get_global_id(1);

    int i = x + y * w;

    if (x >= w || y >= h) return;

    float a0 = alpha[i];
    float a1 = 1.0f - a0;

    float b00 = a0 * image[i * 3 + 0];
    float b01 = a0 * image[i * 3 + 1];
    float b02 = a0 * image[i * 3 + 2];

    float b10 = a1 * image[i * 3 + 0];
    float b11 = a1 * image[i * 3 + 1];
    float b12 = a1 * image[i * 3 + 2];

    int js[4] = {
        max(    0, x - 1) + y * w,
        min(w - 1, x + 1) + y * w,
        x + max(    0, y - 1) * w,
        x + min(h - 1, y + 1) * w,
    };

    float a_sum = 0.0f;

    for (int d = 0; d < 4; d++){
        int j = js[d];

        float da = regularization + fabs(a0 - alpha[j]);

        a_sum += da;

        b00 += da * F_prev[j * 3 + 0];
        b01 += da * F_prev[j * 3 + 1];
        b02 += da * F_prev[j * 3 + 2];

        b10 += da * B_prev[j * 3 + 0];
        b11 += da * B_prev[j * 3 + 1];
        b12 += da * B_prev[j * 3 + 2];
    }

    float a00 = a0 * a0 + a_sum;
    float a11 = a1 * a1 + a_sum;
    float a01 = a0 * a1;

    float inv_det = 1.0f / (a00 * a11 - a01 * a01);

    F[i * 3 + 0] = fmax(0.0f, fmin(1.0f, inv_det * (a11 * b00 - a01 * b10)));
    F[i * 3 + 1] = fmax(0.0f, fmin(1.0f, inv_det * (a11 * b01 - a01 * b11)));
    F[i * 3 + 2] = fmax(0.0f, fmin(1.0f, inv_det * (a11 * b02 - a01 * b12)));
    B[i * 3 + 0] = fmax(0.0f, fmin(1.0f, inv_det * (a00 * b10 - a01 * b00)));
    B[i * 3 + 1] = fmax(0.0f, fmin(1.0f, inv_det * (a00 * b11 - a01 * b01)));
    B[i * 3 + 2] = fmax(0.0f, fmin(1.0f, inv_det * (a00 * b12 - a01 * b02)));
}
h㈵>
          Fc                  C   s  dd }dd }dd }	| j \}
}}|dksJ |
| | }|| } ||}||}||}||}||}||}||
| }t||
d  }d	d
 }d}d}||| ||
||| ||| ||
||| t|d D ]u}t|||  }t|
||  }||| ||
||| |||||
||d |||||||| |||||||| |}t|||kr|}t|D ]'}tt||fd||||||t	
|t	
|t	| ||}}||}}q|}|}qk|	||
||f}|	||
||f}||||| |||fD ]}|  q|r	||fS |S )z=See the :code:`estimate_foreground` method for documentation.c                 S   s,   |  tj }tjttjjtjj	B |dS )N)hostbuf)
astypenpfloat32flattenclBuffercontext	mem_flags	READ_ONLYCOPY_HOST_PTR)arrayr    r   [/tmp/pip-target-vg8gfxp4/lib/python/pymatting/foreground/estimate_foreground_ml_pyopencl.pyuploads   s   z/estimate_foreground_ml_pyopencl.<locals>.uploadc                  W   s    t | }tttjj|d S )N   )r   productr   r   r   r   
READ_WRITE)shapenr   r   r   alloc{   s   
z.estimate_foreground_ml_pyopencl.<locals>.allocc                 S   s(   t j|t jd}tt||  ||S )N)dtype)r   emptyr	   r   enqueue_copyqueuereshape)
device_bufr   host_bufr   r   r   download   s   
z1estimate_foreground_ml_pyopencl.<locals>.download      c              	   S   s2   t jt||fd | |gt|||||gR   d S )N)programresize_nearestr   r   int32)dstsrcw_srch_srcw_dsth_dstdepthr   r   r   r&      s   z7estimate_foreground_ml_pyopencl.<locals>.resize_nearestN)r   max
bit_lengthrangeroundminr%   ml_iterationr   r   r'   r	   release) input_imageinput_alpharegularizationn_small_iterationsn_big_iterations
small_sizereturn_backgroundr   r   r"   h0w0r.   r   F_prevB_prevFBimage_levelalpha_leveln_levelsr&   w_prevh_previ_levelwhn_iteri_iterF_hostB_hostbufr   r   r   estimate_foreground_ml_pyopenclh   s   



rP   )r   r   r   r   F)numpyr   pyopenclr   sourceget_platformsplatformget_devicesdevice_typeGPUdeviceContextr   CommandQueuer   Programbuildr%   rP   r   r   r   r   <module>   s    ]
