smooth_filter.py

"""
Copyright (C) 2018 NVIDIA Corporation.  All rights reserved.
Licensed under the CC BY-NC-SA 4.0 license (https://creativecommons.org/licenses/by-nc-sa/4.0/legalcode).
"""
src = '''
	#include "/usr/local/cuda/include/math_functions.h"
	#define TB 256
	#define EPS 1e-7

	__device__ bool InverseMat4x4(double m_in[4][4], double inv_out[4][4]) {
		double m[16], inv[16];
		for (int i = 0; i < 4; i++) {
			for (int j = 0; j < 4; j++) {
				m[i * 4 + j] = m_in[i][j];
			}
		}

	    inv[0] = m[5]  * m[10] * m[15] -
	             m[5]  * m[11] * m[14] -
	             m[9]  * m[6]  * m[15] +
	             m[9]  * m[7]  * m[14] +
	             m[13] * m[6]  * m[11] -
	             m[13] * m[7]  * m[10];

	    inv[4] = -m[4]  * m[10] * m[15] +
	              m[4]  * m[11] * m[14] +
	              m[8]  * m[6]  * m[15] -
	              m[8]  * m[7]  * m[14] -
	              m[12] * m[6]  * m[11] +
	              m[12] * m[7]  * m[10];

	    inv[8] = m[4]  * m[9] * m[15] -
	             m[4]  * m[11] * m[13] -
	             m[8]  * m[5] * m[15] +
	             m[8]  * m[7] * m[13] +
	             m[12] * m[5] * m[11] -
	             m[12] * m[7] * m[9];

	    inv[12] = -m[4]  * m[9] * m[14] +
	               m[4]  * m[10] * m[13] +
	               m[8]  * m[5] * m[14] -
	               m[8]  * m[6] * m[13] -
	               m[12] * m[5] * m[10] +
	               m[12] * m[6] * m[9];

	    inv[1] = -m[1]  * m[10] * m[15] +
	              m[1]  * m[11] * m[14] +
	              m[9]  * m[2] * m[15] -
	              m[9]  * m[3] * m[14] -
	              m[13] * m[2] * m[11] +
	              m[13] * m[3] * m[10];

	    inv[5] = m[0]  * m[10] * m[15] -
	             m[0]  * m[11] * m[14] -
	             m[8]  * m[2] * m[15] +
	             m[8]  * m[3] * m[14] +
	             m[12] * m[2] * m[11] -
	             m[12] * m[3] * m[10];

	    inv[9] = -m[0]  * m[9] * m[15] +
	              m[0]  * m[11] * m[13] +
	              m[8]  * m[1] * m[15] -
	              m[8]  * m[3] * m[13] -
	              m[12] * m[1] * m[11] +
	              m[12] * m[3] * m[9];

	    inv[13] = m[0]  * m[9] * m[14] -
	              m[0]  * m[10] * m[13] -
	              m[8]  * m[1] * m[14] +
	              m[8]  * m[2] * m[13] +
	              m[12] * m[1] * m[10] -
	              m[12] * m[2] * m[9];

	    inv[2] = m[1]  * m[6] * m[15] -
	             m[1]  * m[7] * m[14] -
	             m[5]  * m[2] * m[15] +
	             m[5]  * m[3] * m[14] +
	             m[13] * m[2] * m[7] -
	             m[13] * m[3] * m[6];

	    inv[6] = -m[0]  * m[6] * m[15] +
	              m[0]  * m[7] * m[14] +
	              m[4]  * m[2] * m[15] -
	              m[4]  * m[3] * m[14] -
	              m[12] * m[2] * m[7] +
	              m[12] * m[3] * m[6];

	    inv[10] = m[0]  * m[5] * m[15] -
	              m[0]  * m[7] * m[13] -
	              m[4]  * m[1] * m[15] +
	              m[4]  * m[3] * m[13] +
	              m[12] * m[1] * m[7] -
	              m[12] * m[3] * m[5];

	    inv[14] = -m[0]  * m[5] * m[14] +
	               m[0]  * m[6] * m[13] +
	               m[4]  * m[1] * m[14] -
	               m[4]  * m[2] * m[13] -
	               m[12] * m[1] * m[6] +
	               m[12] * m[2] * m[5];

	    inv[3] = -m[1] * m[6] * m[11] +
	              m[1] * m[7] * m[10] +
	              m[5] * m[2] * m[11] -
	              m[5] * m[3] * m[10] -
	              m[9] * m[2] * m[7] +
	              m[9] * m[3] * m[6];

	    inv[7] = m[0] * m[6] * m[11] -
	             m[0] * m[7] * m[10] -
	             m[4] * m[2] * m[11] +
	             m[4] * m[3] * m[10] +
	             m[8] * m[2] * m[7] -
	             m[8] * m[3] * m[6];

	    inv[11] = -m[0] * m[5] * m[11] +
	               m[0] * m[7] * m[9] +
	               m[4] * m[1] * m[11] -
	               m[4] * m[3] * m[9] -
	               m[8] * m[1] * m[7] +
	               m[8] * m[3] * m[5];

	    inv[15] = m[0] * m[5] * m[10] -
	              m[0] * m[6] * m[9] -
	              m[4] * m[1] * m[10] +
	              m[4] * m[2] * m[9] +
	              m[8] * m[1] * m[6] -
	              m[8] * m[2] * m[5];

	    double det = m[0] * inv[0] + m[1] * inv[4] + m[2] * inv[8] + m[3] * inv[12];

	    if (abs(det) < 1e-9) {
	        return false;
	    }


	    det = 1.0 / det;

	    for (int i = 0; i < 4; i++) {
	    	for (int j = 0; j < 4; j++) {
	    		inv_out[i][j] = inv[i * 4 + j] * det;
	    	}
	    }

	    return true;
	}

  extern "C"
	__global__ void best_local_affine_kernel(
		float *output, float *input, float *affine_model,
		int h, int w, float epsilon, int kernel_radius
	)
	{
		int size = h * w;
		int id = blockIdx.x * blockDim.x + threadIdx.x;

		if (id < size) {
			int x = id % w, y = id / w;

			double Mt_M[4][4] = {}; // 4x4
			double invMt_M[4][4] = {};
			double Mt_S[3][4] = {}; // RGB -> 1x4
			double A[3][4] = {};
			for (int i = 0; i < 4; i++)
				for (int j = 0; j < 4; j++) {
					Mt_M[i][j] = 0, invMt_M[i][j] = 0;
					if (i != 3) {
						Mt_S[i][j] = 0, A[i][j] = 0;
						if (i == j)
				    		Mt_M[i][j] = 1e-3;
				    }
				}

			for (int dy = -kernel_radius; dy <= kernel_radius; dy++) {
				for (int dx = -kernel_radius; dx <= kernel_radius; dx++) {

					int xx = x + dx, yy = y + dy;
					int id2 = yy * w + xx;

					if (0 <= xx && xx < w && 0 <= yy && yy < h) {

						Mt_M[0][0] += input[id2 + 2*size] * input[id2 + 2*size];
						Mt_M[0][1] += input[id2 + 2*size] * input[id2 + size];
						Mt_M[0][2] += input[id2 + 2*size] * input[id2];
						Mt_M[0][3] += input[id2 + 2*size];

						Mt_M[1][0] += input[id2 + size] * input[id2 + 2*size];
						Mt_M[1][1] += input[id2 + size] * input[id2 + size];
						Mt_M[1][2] += input[id2 + size] * input[id2];
						Mt_M[1][3] += input[id2 + size];

						Mt_M[2][0] += input[id2] * input[id2 + 2*size];
						Mt_M[2][1] += input[id2] * input[id2 + size];
						Mt_M[2][2] += input[id2] * input[id2];
						Mt_M[2][3] += input[id2];

						Mt_M[3][0] += input[id2 + 2*size];
						Mt_M[3][1] += input[id2 + size];
						Mt_M[3][2] += input[id2];
						Mt_M[3][3] += 1;

						Mt_S[0][0] += input[id2 + 2*size] * output[id2 + 2*size];
						Mt_S[0][1] += input[id2 + size] * output[id2 + 2*size];
						Mt_S[0][2] += input[id2] * output[id2 + 2*size];
						Mt_S[0][3] += output[id2 + 2*size];

						Mt_S[1][0] += input[id2 + 2*size] * output[id2 + size];
						Mt_S[1][1] += input[id2 + size] * output[id2 + size];
						Mt_S[1][2] += input[id2] * output[id2 + size];
						Mt_S[1][3] += output[id2 + size];

						Mt_S[2][0] += input[id2 + 2*size] * output[id2];
						Mt_S[2][1] += input[id2 + size] * output[id2];
						Mt_S[2][2] += input[id2] * output[id2];
						Mt_S[2][3] += output[id2];
					}
				}
			}

			bool success = InverseMat4x4(Mt_M, invMt_M);

			for (int i = 0; i < 3; i++) {
				for (int j = 0; j < 4; j++) {
					for (int k = 0; k < 4; k++) {
						A[i][j] += invMt_M[j][k] * Mt_S[i][k];
					}
				}
			}

			for (int i = 0; i < 3; i++) {
				for (int j = 0; j < 4; j++) {
					int affine_id = i * 4 + j;
					affine_model[12 * id + affine_id] = A[i][j];
				}
			}
		}
		return ;
	}

  extern "C"
	__global__ void bilateral_smooth_kernel(
		float *affine_model, float *filtered_affine_model, float *guide,
		int h, int w, int kernel_radius, float sigma1, float sigma2
	)
	{
		int id = blockIdx.x * blockDim.x + threadIdx.x;
		int size = h * w;
		if (id < size) {
			int x = id % w;
			int y = id / w;

			double sum_affine[12] = {};
			double sum_weight = 0;
			for (int dx = -kernel_radius; dx <= kernel_radius; dx++) {
				for (int dy = -kernel_radius; dy <= kernel_radius; dy++) {
					int yy = y + dy, xx = x + dx;
					int id2 = yy * w + xx;
					if (0 <= xx && xx < w && 0 <= yy && yy < h) {
						float color_diff1 = guide[yy*w + xx] - guide[y*w + x];
						float color_diff2 = guide[yy*w + xx + size] - guide[y*w + x + size];
						float color_diff3 = guide[yy*w + xx + 2*size] - guide[y*w + x + 2*size];
						float color_diff_sqr =
							(color_diff1*color_diff1 + color_diff2*color_diff2 + color_diff3*color_diff3) / 3;

						float v1 = exp(-(dx * dx + dy * dy) / (2 * sigma1 * sigma1));
						float v2 = exp(-(color_diff_sqr) / (2 * sigma2 * sigma2));
						float weight = v1 * v2;

						for (int i = 0; i < 3; i++) {
							for (int j = 0; j < 4; j++) {
								int affine_id = i * 4 + j;
								sum_affine[affine_id] += weight * affine_model[id2*12 + affine_id];
							}
						}
						sum_weight += weight;
					}
				}
			}

			for (int i = 0; i < 3; i++) {
				for (int j = 0; j < 4; j++) {
					int affine_id = i * 4 + j;
					filtered_affine_model[id*12 + affine_id] = sum_affine[affine_id] / sum_weight;
				}
			}
		}
		return ;
	}


  extern "C"
	__global__ void reconstruction_best_kernel(
		float *input, float *filtered_affine_model, float *filtered_best_output,
		int h, int w
	)
	{
		int id = blockIdx.x * blockDim.x + threadIdx.x;
		int size = h * w;
		if (id < size) {
			double out1 =
				input[id + 2*size] * filtered_affine_model[id*12 + 0] + // A[0][0] +
				input[id + size]   * filtered_affine_model[id*12 + 1] + // A[0][1] +
				input[id]          * filtered_affine_model[id*12 + 2] + // A[0][2] +
									 filtered_affine_model[id*12 + 3]; //A[0][3];
			double out2 =
				input[id + 2*size] * filtered_affine_model[id*12 + 4] + //A[1][0] +
				input[id + size]   * filtered_affine_model[id*12 + 5] + //A[1][1] +
				input[id]          * filtered_affine_model[id*12 + 6] + //A[1][2] +
									 filtered_affine_model[id*12 + 7]; //A[1][3];
			double out3 =
				input[id + 2*size] * filtered_affine_model[id*12 + 8] + //A[2][0] +
				input[id + size]   * filtered_affine_model[id*12 + 9] + //A[2][1] +
				input[id]          * filtered_affine_model[id*12 + 10] + //A[2][2] +
									 filtered_affine_model[id*12 + 11]; // A[2][3];

			filtered_best_output[id] = out1;
			filtered_best_output[id + size] = out2;
			filtered_best_output[id + 2*size] = out3;
		}
		return ;
	}
	'''

import torch
import numpy as np
from PIL import Image
from cupy.cuda import function
from pynvrtc.compiler import Program
from collections import namedtuple


def smooth_local_affine(output_cpu, input_cpu, epsilon, patch, h, w, f_r, f_e):
    # program = Program(src.encode('utf-8'), 'best_local_affine_kernel.cu'.encode('utf-8'))
    # ptx = program.compile(['-I/usr/local/cuda/include'.encode('utf-8')])
    program = Program(src, 'best_local_affine_kernel.cu')
    ptx = program.compile(['-I/usr/local/cuda/include'])
    m = function.Module()
    m.load(bytes(ptx.encode()))

    _reconstruction_best_kernel = m.get_function('reconstruction_best_kernel')
    _bilateral_smooth_kernel = m.get_function('bilateral_smooth_kernel')
    _best_local_affine_kernel = m.get_function('best_local_affine_kernel')
    Stream = namedtuple('Stream', ['ptr'])
    s = Stream(ptr=torch.cuda.current_stream().cuda_stream)

    filter_radius = f_r
    sigma1 = filter_radius / 3
    sigma2 = f_e
    radius = (patch - 1) / 2

    filtered_best_output = torch.zeros(np.shape(input_cpu)).cuda()
    affine_model =    torch.zeros((h * w, 12)).cuda()
    filtered_affine_model =torch.zeros((h * w, 12)).cuda()

    input_ = torch.from_numpy(input_cpu).cuda()
    output_ = torch.from_numpy(output_cpu).cuda()
    _best_local_affine_kernel(
        grid=(int((h * w) / 256 + 1), 1),
        block=(256, 1, 1),
        args=[output_.data_ptr(), input_.data_ptr(), affine_model.data_ptr(),
             np.int32(h), np.int32(w), np.float32(epsilon), np.int32(radius)], stream=s
     )

    _bilateral_smooth_kernel(
        grid=(int((h * w) / 256 + 1), 1),
        block=(256, 1, 1),
        args=[affine_model.data_ptr(), filtered_affine_model.data_ptr(), input_.data_ptr(), np.int32(h), np.int32(w), np.int32(f_r), np.float32(sigma1), np.float32(sigma2)], stream=s
    )

    _reconstruction_best_kernel(
        grid=(int((h * w) / 256 + 1), 1),
        block=(256, 1, 1),
        args=[input_.data_ptr(), filtered_affine_model.data_ptr(), filtered_best_output.data_ptr(),
        np.int32(h), np.int32(w)], stream=s
    )
    numpy_filtered_best_output = filtered_best_output.cpu().numpy()
    return numpy_filtered_best_output


def smooth_filter(initImg, contentImg, f_radius=15,f_edge=1e-1):
    '''
    :param initImg: intermediate output. Either image path or PIL Image
    :param contentImg: content image output. Either path or PIL Image
    :return: stylized output image. PIL Image
    '''
    if type(initImg) == str:
        initImg = Image.open(initImg).convert("RGB")
    best_image_bgr = np.array(initImg, dtype=np.float32)
    bW, bH, bC = best_image_bgr.shape
    best_image_bgr = best_image_bgr[:, :, ::-1]
    best_image_bgr = best_image_bgr.transpose((2, 0, 1))

    if type(contentImg) == str:
        contentImg = Image.open(contentImg).convert("RGB")
    content_input = contentImg.resize((bH,bW))
    content_input = np.array(content_input, dtype=np.float32)
    content_input = content_input[:, :, ::-1]
    content_input = content_input.transpose((2, 0, 1))
    input_ = np.ascontiguousarray(content_input, dtype=np.float32) / 255.
    _, H, W = np.shape(input_)
    output_ = np.ascontiguousarray(best_image_bgr, dtype=np.float32) / 255.
    best_ = smooth_local_affine(output_, input_, 1e-7, 3, H, W, f_radius, f_edge)
    best_ = best_.transpose(1, 2, 0)
    result = Image.fromarray(np.uint8(np.clip(best_ * 255., 0, 255.)))
    return result