Function buildLinearSystem

src/backend/cuda/kernel/homography.hpp:200–295

template<typename T>
__global__ void buildLinearSystem(Param<T> H, CParam<float> x_src,
                                  CParam<float> y_src, CParam<float> x_dst,
                                  CParam<float> y_dst, CParam<float> rnd,
                                  const unsigned iterations) {
    unsigned tid_y = threadIdx.y;
    unsigned i     = blockIdx.y * blockDim.y + tid_y;

    if (i < iterations) {
        float x_src_mean, y_src_mean;
        float x_dst_mean, y_dst_mean;
        float src_scale, dst_scale;
        float src_pt_x[4], src_pt_y[4], dst_pt_x[4], dst_pt_y[4];

        computeMeanScale(&x_src_mean, &y_src_mean, &x_dst_mean, &y_dst_mean,
                         &src_scale, &dst_scale, src_pt_x, src_pt_y, dst_pt_x,
                         dst_pt_y, x_src, y_src, x_dst, y_dst, rnd, i);

        T* s_V = (T*)sh;
        T* s_S = (T*)sh + 16 * 81;

        // Compute input matrix
        for (unsigned j = threadIdx.x; j < 4; j += blockDim.x) {
            float srcx = (src_pt_x[j] - x_src_mean) * src_scale;
            float srcy = (src_pt_y[j] - y_src_mean) * src_scale;
            float dstx = (dst_pt_x[j] - x_dst_mean) * dst_scale;
            float dsty = (dst_pt_y[j] - y_dst_mean) * dst_scale;

            SSPTR(tid_y, 0, j * 2) = 0.0f;
            SSPTR(tid_y, 1, j * 2) = 0.0f;
            SSPTR(tid_y, 2, j * 2) = 0.0f;
            SSPTR(tid_y, 3, j * 2) = -srcx;
            SSPTR(tid_y, 4, j * 2) = -srcy;
            SSPTR(tid_y, 5, j * 2) = -1.0f;
            SSPTR(tid_y, 6, j * 2) = dsty * srcx;
            SSPTR(tid_y, 7, j * 2) = dsty * srcy;
            SSPTR(tid_y, 8, j * 2) = dsty;

            SSPTR(tid_y, 0, j * 2 + 1) = srcx;
            SSPTR(tid_y, 1, j * 2 + 1) = srcy;
            SSPTR(tid_y, 2, j * 2 + 1) = 1.0f;
            SSPTR(tid_y, 3, j * 2 + 1) = 0.0f;
            SSPTR(tid_y, 4, j * 2 + 1) = 0.0f;
            SSPTR(tid_y, 5, j * 2 + 1) = 0.0f;
            SSPTR(tid_y, 6, j * 2 + 1) = -dstx * srcx;
            SSPTR(tid_y, 7, j * 2 + 1) = -dstx * srcy;
            SSPTR(tid_y, 8, j * 2 + 1) = -dstx;

            if (j == 4) {
                SSPTR(tid_y, 0, 8) = 0.0f;
                SSPTR(tid_y, 1, 8) = 0.0f;
                SSPTR(tid_y, 2, 8) = 0.0f;
                SSPTR(tid_y, 3, 8) = 0.0f;
                SSPTR(tid_y, 4, 8) = 0.0f;
                SSPTR(tid_y, 5, 8) = 0.0f;
                SSPTR(tid_y, 6, 8) = 0.0f;
                SSPTR(tid_y, 7, 8) = 0.0f;
                SSPTR(tid_y, 8, 8) = 0.0f;
            }